975 lines · plain
1; This is an excerpt from the SYCL end-to-end test suite, cleaned out from unrelevant details,2; that reproduced multiple cases of the issues when OpPhi's result type mismatches with operand types.3; The only pass criterion is that spirv-val considers output valid.4 5; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val --target-env spv1.4 %}6 7%struct.PFWGFunctor = type { i64, i64, i32, i32, %"class.sycl::_V1::accessor" }8%"class.sycl::_V1::accessor" = type { %"class.sycl::_V1::detail::AccessorImplDevice", %union.anon }9%"class.sycl::_V1::detail::AccessorImplDevice" = type { %"class.sycl::_V1::range", %"class.sycl::_V1::range", %"class.sycl::_V1::range" }10%"class.sycl::_V1::range" = type { %"class.sycl::_V1::detail::array" }11%"class.sycl::_V1::detail::array" = type { [1 x i64] }12%union.anon = type { ptr addrspace(1) }13%class.anon.2 = type { %"class.sycl::_V1::accessor" }14%"class.sycl::_V1::group" = type { %"class.sycl::_V1::range", %"class.sycl::_V1::range", %"class.sycl::_V1::range", %"class.sycl::_V1::range" }15%"class.sycl::_V1::group.15" = type { %"class.sycl::_V1::range.16", %"class.sycl::_V1::range.16", %"class.sycl::_V1::range.16", %"class.sycl::_V1::range.16" }16%"class.sycl::_V1::range.16" = type { %"class.sycl::_V1::detail::array.17" }17%"class.sycl::_V1::detail::array.17" = type { [2 x i64] }18%"class.sycl::_V1::private_memory" = type { %struct.MyStruct }19%struct.MyStruct = type { i32, i32 }20 21@GFunctor = internal addrspace(3) global %struct.PFWGFunctor undef, align 822@WI.0 = internal unnamed_addr addrspace(3) global i64 undef, align 823@WI.1 = internal unnamed_addr addrspace(3) global i64 undef, align 824@WI.2 = internal unnamed_addr addrspace(3) global i64 undef, align 825@WI.3 = internal unnamed_addr addrspace(3) global i64 undef, align 826@WI.4 = internal unnamed_addr addrspace(3) global i32 undef, align 827@WI.6 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 828@GCnt = internal unnamed_addr addrspace(3) global i32 undef, align 429@__spirv_BuiltInNumWorkgroups = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 3230@GKernel1 = internal addrspace(3) global %class.anon.2 undef, align 831@GCnt2 = internal unnamed_addr addrspace(3) global i32 undef, align 432@GKernel2 = internal addrspace(3) global %class.anon.2 undef, align 833@GCnt3 = internal unnamed_addr addrspace(3) global i32 undef, align 434@GKernel3 = internal addrspace(3) global %class.anon.2 undef, align 835@GCnt4 = internal unnamed_addr addrspace(3) global i32 undef, align 436@GKernel4 = internal addrspace(3) global %class.anon.2 undef, align 837@GCnt5 = internal unnamed_addr addrspace(3) global i32 undef, align 438@__spirv_BuiltInLocalInvocationIndex = external local_unnamed_addr addrspace(1) constant i64, align 839@GThis = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 840@GAsCast = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 841@GCmp = internal unnamed_addr addrspace(3) global i1 undef, align 142@WGCopy = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 843@WGCopy.1.0 = internal unnamed_addr addrspace(3) global i64 undef, align 1644@WGCopy.1.1 = internal unnamed_addr addrspace(3) global i64 undef, align 1645@WGCopy.1.2 = internal unnamed_addr addrspace(3) global i64 undef, align 1646@WGCopy.1.3 = internal unnamed_addr addrspace(3) global i64 undef, align 1647@WGCopy.1.4 = internal unnamed_addr addrspace(3) global i32 undef, align 1648@WGCopy.1.5 = internal unnamed_addr addrspace(3) global i32 undef, align 1649@WGCopy.1.6 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 1650@ArgShadow = internal unnamed_addr addrspace(3) global %"class.sycl::_V1::group" undef, align 1651@GAsCast2 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 852@GCmp2 = internal unnamed_addr addrspace(3) global i1 undef, align 153@WGCopy.3.0 = internal unnamed_addr addrspace(3) global i64 undef, align 854@WGCopy.4.0 = internal unnamed_addr addrspace(3) global i64 undef, align 855@WGCopy.5.0 = internal unnamed_addr addrspace(3) global i64 undef, align 856@WGCopy.6.0 = internal unnamed_addr addrspace(3) global i64 undef, align 857@ArgShadow.7 = internal unnamed_addr addrspace(3) global %"class.sycl::_V1::group" undef, align 1658@GAscast3 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 859@GCmp3 = internal unnamed_addr addrspace(3) global i1 undef, align 160@WGCopy.9.0 = internal unnamed_addr addrspace(3) global i64 undef, align 861@WGCopy.10.0 = internal unnamed_addr addrspace(3) global i64 undef, align 862@ArgShadow.11 = internal unnamed_addr addrspace(3) global %"class.sycl::_V1::group" undef, align 1663@GAsCast4 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 864@GCmp4 = internal unnamed_addr addrspace(3) global i1 undef, align 165@WGCopy.13.0 = internal unnamed_addr addrspace(3) global i64 undef, align 866@WGCopy.13.1 = internal unnamed_addr addrspace(3) global i64 undef, align 867@WGCopy.14.0 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 868@WGCopy.14.1 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 869@WGCopy.15.0 = internal unnamed_addr addrspace(3) global i64 undef, align 870@WGCopy.15.1 = internal unnamed_addr addrspace(3) global i64 undef, align 871@WGCopy.16.0 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 872@WGCopy.16.1 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 873@ArgShadow.17 = internal unnamed_addr addrspace(3) global %"class.sycl::_V1::group.15" undef, align 1674@GAsCast5 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 875@GCmp5 = internal unnamed_addr addrspace(3) global i1 undef, align 176@WGCopy.19.0 = internal unnamed_addr addrspace(3) global i64 undef, align 877@WGCopy.20.0 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 878@WGCopy.20.1 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 879@ArgShadow.21 = internal unnamed_addr addrspace(3) global %"class.sycl::_V1::group" undef, align 1680@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 3281@__spirv_BuiltInGlobalSize = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 3282@__spirv_BuiltInLocalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 3283@__spirv_BuiltInWorkgroupId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 3284@__spirv_BuiltInWorkgroupSize = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 3285 86; Function Attrs: convergent mustprogress norecurse nounwind87define weak_odr dso_local spir_kernel void @_ZTS11PFWGFunctor(i64 noundef %_arg_wg_chunk, i64 noundef %_arg_range_length, i32 noundef %_arg_n_iter, i32 noundef %_arg_addend, ptr addrspace(1) noundef align 4 %_arg_dev_ptr, ptr noundef byval(%"class.sycl::_V1::range") align 8 %_arg_dev_ptr1, ptr noundef byval(%"class.sycl::_V1::range") align 8 %_arg_dev_ptr2, ptr noundef byval(%"class.sycl::_V1::range") align 8 %_arg_dev_ptr3) {88entry:89 %agg.tmp67 = alloca %"class.sycl::_V1::group", align 890 store i64 %_arg_wg_chunk, ptr addrspace(3) @GFunctor, align 891 store i64 %_arg_range_length, ptr addrspace(3) undef, align 892 store i32 %_arg_n_iter, ptr addrspace(3) undef, align 893 store i32 %_arg_addend, ptr addrspace(3) undef, align 494 %0 = load i64, ptr %_arg_dev_ptr1, align 895 %1 = load i64, ptr %_arg_dev_ptr2, align 896 %2 = load i64, ptr %_arg_dev_ptr3, align 897 store i64 %2, ptr addrspace(3) undef, align 898 store i64 %0, ptr addrspace(3) undef, align 899 store i64 %1, ptr addrspace(3) undef, align 8100 %add.ptr.i = getelementptr inbounds i32, ptr addrspace(1) %_arg_dev_ptr, i64 %2101 store ptr addrspace(1) %add.ptr.i, ptr addrspace(3) undef, align 8102 %3 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalSize, align 32103 %4 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupSize, align 32104 %5 = load i64, ptr addrspace(1) @__spirv_BuiltInNumWorkgroups, align 32105 %6 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupId, align 32106 call void @llvm.lifetime.start.p0(i64 32, ptr nonnull %agg.tmp67)107 store i64 %3, ptr %agg.tmp67, align 1108 %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 8109 store i64 %4, ptr %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx, align 1110 %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 16111 store i64 %5, ptr %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx, align 1112 %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 24113 store i64 %6, ptr %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx, align 1114 %7 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationIndex, align 8115 tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)116 %cmpz15.i = icmp eq i64 %7, 0117 br i1 %cmpz15.i, label %leader.i, label %merge.i118 119leader.i: ; preds = %entry120 call void @llvm.memcpy.p3.p0.i64(ptr addrspace(3) noundef align 16 dereferenceable(32) @ArgShadow, ptr noundef nonnull align 8 dereferenceable(32) %agg.tmp67, i64 32, i1 false)121 br label %merge.i122 123merge.i: ; preds = %leader.i, %entry124 tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)125 call void @llvm.memcpy.p0.p3.i64(ptr noundef nonnull align 8 dereferenceable(32) %agg.tmp67, ptr addrspace(3) noundef align 16 dereferenceable(32) @ArgShadow, i64 32, i1 false)126 tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)127 br i1 %cmpz15.i, label %wg_leader.i, label %wg_cf.i128 129wg_leader.i: ; preds = %merge.i130 %g.ascast.i = addrspacecast ptr %agg.tmp67 to ptr addrspace(4)131 store ptr addrspace(4) %g.ascast.i, ptr addrspace(3) @GAsCast, align 8132 store ptr addrspace(4) addrspacecast (ptr addrspace(3) @GFunctor to ptr addrspace(4)), ptr addrspace(3) @GThis, align 8133 %8 = load i32, ptr addrspace(3) undef, align 4134 %9 = load i64, ptr addrspace(3) @GFunctor, align 8135 %index.i = getelementptr inbounds i8, ptr %agg.tmp67, i64 24136 %10 = load i64, ptr %index.i, align 8137 %mul.i = mul i64 %9, %10138 %localRange.i = getelementptr inbounds i8, ptr %agg.tmp67, i64 8139 %11 = load i64, ptr %localRange.i, align 8140 %12 = load i64, ptr addrspace(3) undef, align 8141 store i64 %9, ptr addrspace(3) @WI.0, align 8142 store i64 %11, ptr addrspace(3) @WI.1, align 8143 store i64 %mul.i, ptr addrspace(3) @WI.2, align 8144 store i64 %12, ptr addrspace(3) @WI.3, align 8145 store i32 %8, ptr addrspace(3) @WI.4, align 8146 store ptr addrspace(4) undef, ptr addrspace(3) @WI.6, align 8147 store i32 0, ptr addrspace(3) @GCnt, align 4148 br label %wg_cf.i149 150wg_cf.i: ; preds = %wg_leader.i, %merge.i151 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)152 %wg_val_this1.i = load ptr addrspace(4), ptr addrspace(3) @GThis, align 8153 %n_iter.i = getelementptr inbounds i8, ptr addrspace(4) %wg_val_this1.i, i64 16154 %13 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationId, align 32155 br label %for.cond.i156 157for.cond.i: ; preds = %wg_cf11.i, %wg_cf.i158 %agg.tmp.i.sroa.0.0 = phi i64 [ undef, %wg_cf.i ], [ %agg.tmp.i.sroa.0.0.copyload13, %wg_cf11.i ]159 %agg.tmp.i.sroa.6.0 = phi i64 [ undef, %wg_cf.i ], [ %agg.tmp.i.sroa.6.0.copyload15, %wg_cf11.i ]160 %agg.tmp.i.sroa.7.0 = phi i64 [ undef, %wg_cf.i ], [ %agg.tmp.i.sroa.7.0.copyload17, %wg_cf11.i ]161 %agg.tmp.i.sroa.8.0 = phi i64 [ undef, %wg_cf.i ], [ %agg.tmp.i.sroa.8.0.copyload19, %wg_cf11.i ]162 %agg.tmp.i.sroa.9.0 = phi i32 [ undef, %wg_cf.i ], [ %agg.tmp.i.sroa.9.0.copyload21, %wg_cf11.i ]163 %agg.tmp.i.sroa.10.0 = phi i32 [ undef, %wg_cf.i ], [ %agg.tmp.i.sroa.10.0.copyload23, %wg_cf11.i ]164 %agg.tmp.i.sroa.11.0 = phi ptr addrspace(4) [ undef, %wg_cf.i ], [ %agg.tmp.i.sroa.11.0.copyload25, %wg_cf11.i ]165 %this.addr.0.i = phi ptr addrspace(4) [ addrspacecast (ptr addrspace(3) @GFunctor to ptr addrspace(4)), %wg_cf.i ], [ %mat_ld13.i, %wg_cf11.i ]166 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)167 br i1 %cmpz15.i, label %wg_leader4.i, label %wg_cf5.i168 169wg_leader4.i: ; preds = %for.cond.i170 %14 = load i32, ptr addrspace(3) @GCnt, align 4171 %15 = load i32, ptr addrspace(4) %n_iter.i, align 8172 %cmp.i = icmp slt i32 %14, %15173 store i1 %cmp.i, ptr addrspace(3) @GCmp, align 1174 br label %wg_cf5.i175 176wg_cf5.i: ; preds = %wg_leader4.i, %for.cond.i177 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)178 %wg_val_cmp.i = load i1, ptr addrspace(3) @GCmp, align 1179 br i1 %wg_val_cmp.i, label %for.body.i, label %_ZNK11PFWGFunctorclEN4sycl3_V15groupILi1EEE.exit180 181for.body.i: ; preds = %wg_cf5.i182 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)183 br i1 %cmpz15.i, label %wg_leader7.i, label %wg_cf8.i184 185wg_leader7.i: ; preds = %for.body.i186 %agg.tmp.i.sroa.0.0.copyload = load i64, ptr addrspace(3) @WI.0, align 8187 %agg.tmp.i.sroa.6.0.copyload = load i64, ptr addrspace(3) @WI.1, align 8188 %agg.tmp.i.sroa.7.0.copyload = load i64, ptr addrspace(3) @WI.2, align 8189 %agg.tmp.i.sroa.8.0.copyload = load i64, ptr addrspace(3) @WI.3, align 8190 %agg.tmp.i.sroa.9.0.copyload = load i32, ptr addrspace(3) @WI.4, align 8191 %agg.tmp.i.sroa.11.0.copyload = load ptr addrspace(4), ptr addrspace(3) @WI.6, align 8192 br label %wg_cf8.i193 194wg_cf8.i: ; preds = %wg_leader7.i, %for.body.i195 %agg.tmp.i.sroa.0.1 = phi i64 [ %agg.tmp.i.sroa.0.0.copyload, %wg_leader7.i ], [ %agg.tmp.i.sroa.0.0, %for.body.i ]196 %agg.tmp.i.sroa.6.1 = phi i64 [ %agg.tmp.i.sroa.6.0.copyload, %wg_leader7.i ], [ %agg.tmp.i.sroa.6.0, %for.body.i ]197 %agg.tmp.i.sroa.7.1 = phi i64 [ %agg.tmp.i.sroa.7.0.copyload, %wg_leader7.i ], [ %agg.tmp.i.sroa.7.0, %for.body.i ]198 %agg.tmp.i.sroa.8.1 = phi i64 [ %agg.tmp.i.sroa.8.0.copyload, %wg_leader7.i ], [ %agg.tmp.i.sroa.8.0, %for.body.i ]199 %agg.tmp.i.sroa.9.1 = phi i32 [ %agg.tmp.i.sroa.9.0.copyload, %wg_leader7.i ], [ %agg.tmp.i.sroa.9.0, %for.body.i ]200 %agg.tmp.i.sroa.11.1 = phi ptr addrspace(4) [ %agg.tmp.i.sroa.11.0.copyload, %wg_leader7.i ], [ %agg.tmp.i.sroa.11.0, %for.body.i ]201 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)202 br i1 %cmpz15.i, label %TestMat.i, label %LeaderMat.i203 204TestMat.i: ; preds = %wg_cf8.i205 store i64 %agg.tmp.i.sroa.0.1, ptr addrspace(3) @WGCopy.1.0, align 16206 store i64 %agg.tmp.i.sroa.6.1, ptr addrspace(3) @WGCopy.1.1, align 16207 store i64 %agg.tmp.i.sroa.7.1, ptr addrspace(3) @WGCopy.1.2, align 16208 store i64 %agg.tmp.i.sroa.8.1, ptr addrspace(3) @WGCopy.1.3, align 16209 store i32 %agg.tmp.i.sroa.9.1, ptr addrspace(3) @WGCopy.1.4, align 16210 store i32 %agg.tmp.i.sroa.10.0, ptr addrspace(3) @WGCopy.1.5, align 16211 store ptr addrspace(4) %agg.tmp.i.sroa.11.1, ptr addrspace(3) @WGCopy.1.6, align 16212 store ptr addrspace(4) %this.addr.0.i, ptr addrspace(3) @WGCopy, align 8213 br label %LeaderMat.i214 215LeaderMat.i: ; preds = %TestMat.i, %wg_cf8.i216 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)217 %mat_ld13.i = load ptr addrspace(4), ptr addrspace(3) @WGCopy, align 8218 %agg.tmp.i.sroa.0.0.copyload13 = load i64, ptr addrspace(3) @WGCopy.1.0, align 16219 %agg.tmp.i.sroa.6.0.copyload15 = load i64, ptr addrspace(3) @WGCopy.1.1, align 16220 %agg.tmp.i.sroa.7.0.copyload17 = load i64, ptr addrspace(3) @WGCopy.1.2, align 16221 %agg.tmp.i.sroa.8.0.copyload19 = load i64, ptr addrspace(3) @WGCopy.1.3, align 16222 %agg.tmp.i.sroa.9.0.copyload21 = load i32, ptr addrspace(3) @WGCopy.1.4, align 16223 %agg.tmp.i.sroa.10.0.copyload23 = load i32, ptr addrspace(3) @WGCopy.1.5, align 16224 %agg.tmp.i.sroa.11.0.copyload25 = load ptr addrspace(4), ptr addrspace(3) @WGCopy.1.6, align 16225 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)226 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef 2, i32 noundef 2, i32 noundef 272)227 %cmp.not.i.i = icmp ult i64 %13, %agg.tmp.i.sroa.0.0.copyload13228 br i1 %cmp.not.i.i, label %if.end.i.i, label %lexit1229 230if.end.i.i: ; preds = %LeaderMat.i231 %add.i.i = add i64 %agg.tmp.i.sroa.0.0.copyload13, %agg.tmp.i.sroa.6.0.copyload15232 %sub.i.i = add i64 %add.i.i, -1233 %div.i.i = udiv i64 %sub.i.i, %agg.tmp.i.sroa.6.0.copyload15234 %mul.i.i = mul i64 %13, %div.i.i235 %add4.i.i = add i64 %agg.tmp.i.sroa.7.0.copyload17, %mul.i.i236 %add6.i.i = add i64 %add4.i.i, %div.i.i237 %.sroa.speculated.i.i = call i64 @llvm.umin.i64(i64 %agg.tmp.i.sroa.8.0.copyload19, i64 %add6.i.i)238 %16 = getelementptr inbounds i8, ptr addrspace(4) %agg.tmp.i.sroa.11.0.copyload25, i64 24239 br label %for.cond.i.i240 241for.cond.i.i: ; preds = %for.body.i.i, %if.end.i.i242 %ind.0.i.i = phi i64 [ %add4.i.i, %if.end.i.i ], [ %inc.i.i, %for.body.i.i ]243 %cmp8.i.i = icmp ult i64 %ind.0.i.i, %.sroa.speculated.i.i244 br i1 %cmp8.i.i, label %for.body.i.i, label %lexit1245 246for.body.i.i: ; preds = %for.cond.i.i247 %17 = load ptr addrspace(1), ptr addrspace(4) %16, align 8248 %arrayidx.i.i.i = getelementptr inbounds i32, ptr addrspace(1) %17, i64 %ind.0.i.i249 %18 = load i32, ptr addrspace(1) %arrayidx.i.i.i, align 4250 %add10.i.i = add nsw i32 %18, %agg.tmp.i.sroa.9.0.copyload21251 store i32 %add10.i.i, ptr addrspace(1) %arrayidx.i.i.i, align 4252 %inc.i.i = add nuw i64 %ind.0.i.i, 1253 br label %for.cond.i.i254 255lexit1: ; preds = %for.cond.i.i, %LeaderMat.i256 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef 2, i32 noundef 2, i32 noundef 272)257 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)258 br i1 %cmpz15.i, label %wg_leader10.i, label %wg_cf11.i259 260wg_leader10.i: ; preds = %lexit1261 %19 = load i32, ptr addrspace(3) @GCnt, align 4262 %inc.i = add nsw i32 %19, 1263 store i32 %inc.i, ptr addrspace(3) @GCnt, align 4264 br label %wg_cf11.i265 266wg_cf11.i: ; preds = %wg_leader10.i, %lexit1267 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)268 br label %for.cond.i269 270_ZNK11PFWGFunctorclEN4sycl3_V15groupILi1EEE.exit: ; preds = %wg_cf5.i271 call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %agg.tmp67)272 ret void273}274 275; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite)276declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture)277 278; Function Attrs: convergent nounwind279declare dso_local spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef, i32 noundef, i32 noundef)280 281; Function Attrs: nocallback nofree nounwind willreturn memory(argmem: readwrite)282declare void @llvm.memcpy.p3.p0.i64(ptr addrspace(3) noalias nocapture writeonly, ptr noalias nocapture readonly, i64, i1 immarg)283 284; Function Attrs: nocallback nofree nounwind willreturn memory(argmem: readwrite)285declare void @llvm.memcpy.p0.p3.i64(ptr noalias nocapture writeonly, ptr addrspace(3) noalias nocapture readonly, i64, i1 immarg)286 287; Function Attrs: nocallback nofree nosync nounwind speculatable willreturn memory(none)288declare i64 @llvm.umin.i64(i64, i64)289 290; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite)291declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture)292 293; Function Attrs: convergent mustprogress norecurse nounwind294define weak_odr dso_local spir_kernel void @bar(ptr addrspace(1) noundef align 4 %_arg_dev_ptr, ptr noundef byval(%"class.sycl::_V1::range") align 8 %_arg_dev_ptr1, ptr noundef byval(%"class.sycl::_V1::range") align 8 %_arg_dev_ptr2, ptr noundef byval(%"class.sycl::_V1::range") align 8 %_arg_dev_ptr3) {295entry:296 %agg.tmp67 = alloca %"class.sycl::_V1::group", align 8297 %0 = load i64, ptr %_arg_dev_ptr1, align 8298 %1 = load i64, ptr %_arg_dev_ptr2, align 8299 %2 = load i64, ptr %_arg_dev_ptr3, align 8300 store i64 %2, ptr addrspace(3) @GKernel1, align 8301 store i64 %0, ptr addrspace(3) undef, align 8302 store i64 %1, ptr addrspace(3) undef, align 8303 %add.ptr.i = getelementptr inbounds i32, ptr addrspace(1) %_arg_dev_ptr, i64 %2304 store ptr addrspace(1) %add.ptr.i, ptr addrspace(3) undef, align 8305 %3 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalSize, align 32306 %4 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupSize, align 32307 %5 = load i64, ptr addrspace(1) @__spirv_BuiltInNumWorkgroups, align 32308 %6 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupId, align 32309 call void @llvm.lifetime.start.p0(i64 32, ptr nonnull %agg.tmp67)310 store i64 %3, ptr %agg.tmp67, align 1311 %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 8312 store i64 %4, ptr %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx, align 1313 %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 16314 store i64 %5, ptr %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx, align 1315 %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 24316 store i64 %6, ptr %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx, align 1317 %7 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationIndex, align 8318 tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)319 %cmpz27.i = icmp eq i64 %7, 0320 br i1 %cmpz27.i, label %leader.i, label %merge.i321 322leader.i: ; preds = %entry323 call void @llvm.memcpy.p3.p0.i64(ptr addrspace(3) noundef align 16 dereferenceable(32) @ArgShadow.7, ptr noundef nonnull align 8 dereferenceable(32) %agg.tmp67, i64 32, i1 false)324 br label %merge.i325 326merge.i: ; preds = %leader.i, %entry327 tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)328 call void @llvm.memcpy.p0.p3.i64(ptr noundef nonnull align 8 dereferenceable(32) %agg.tmp67, ptr addrspace(3) noundef align 16 dereferenceable(32) @ArgShadow.7, i64 32, i1 false)329 tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)330 br i1 %cmpz27.i, label %wg_leader.i, label %wg_cf.i331 332wg_leader.i: ; preds = %merge.i333 %g.ascast.i = addrspacecast ptr %agg.tmp67 to ptr addrspace(4)334 store ptr addrspace(4) %g.ascast.i, ptr addrspace(3) @GAsCast2, align 8335 store i32 0, ptr addrspace(3) @GCnt2, align 4336 br label %wg_cf.i337 338wg_cf.i: ; preds = %wg_leader.i, %merge.i339 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)340 %8 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalInvocationId, align 32341 %9 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationId, align 32342 %cmp.i.i.i.i.i.i = icmp ult i64 %8, 2147483648343 br label %for.cond.i344 345for.cond.i: ; preds = %wg_cf18.i, %wg_cf.i346 %agg.tmp5.i.sroa.0.0 = phi i64 [ undef, %wg_cf.i ], [ %18, %wg_cf18.i ]347 %agg.tmp4.i.sroa.0.0 = phi i64 [ undef, %wg_cf.i ], [ %17, %wg_cf18.i ]348 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)349 br i1 %cmpz27.i, label %wg_leader8.i, label %wg_cf9.i350 351wg_leader8.i: ; preds = %for.cond.i352 %10 = load i32, ptr addrspace(3) @GCnt2, align 4353 %cmp.i = icmp slt i32 %10, 2354 store i1 %cmp.i, ptr addrspace(3) @GCmp2, align 1355 br label %wg_cf9.i356 357wg_cf9.i: ; preds = %wg_leader8.i, %for.cond.i358 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)359 %wg_val_cmp.i = load i1, ptr addrspace(3) @GCmp2, align 1360 br i1 %wg_val_cmp.i, label %for.body.i, label %lexit2361 362for.body.i: ; preds = %wg_cf9.i363 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)364 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)365 br i1 %cmpz27.i, label %TestMat25.i, label %LeaderMat22.i366 367TestMat25.i: ; preds = %for.body.i368 store i64 %agg.tmp5.i.sroa.0.0, ptr addrspace(3) @WGCopy.6.0, align 8369 store i64 ptrtoint (ptr addrspace(4) addrspacecast (ptr addrspace(3) @GKernel1 to ptr addrspace(4)) to i64), ptr addrspace(3) @WGCopy.4.0, align 8370 store i64 5, ptr addrspace(3) @WGCopy.3.0, align 8371 store i64 %agg.tmp4.i.sroa.0.0, ptr addrspace(3) @WGCopy.5.0, align 8372 br label %LeaderMat22.i373 374LeaderMat22.i: ; preds = %TestMat25.i, %for.body.i375 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)376 %11 = load i64, ptr addrspace(3) @WGCopy.3.0, align 8377 %12 = load i64, ptr addrspace(3) @WGCopy.4.0, align 8378 %13 = inttoptr i64 %12 to ptr addrspace(4)379 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)380 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef 2, i32 noundef 2, i32 noundef 272)381 %14 = getelementptr inbounds i8, ptr addrspace(4) %13, i64 24382 br label %for.cond.i.i383 384for.cond.i.i: ; preds = %for.body.i.i, %LeaderMat22.i385 %storemerge.i.i = phi i64 [ %9, %LeaderMat22.i ], [ %add.i.i, %for.body.i.i ]386 %cmp.i.i = icmp ult i64 %storemerge.i.i, %11387 br i1 %cmp.i.i, label %for.body.i.i, label %lexit3388 389for.body.i.i: ; preds = %for.cond.i.i390 call void @llvm.assume(i1 %cmp.i.i.i.i.i.i)391 %15 = load ptr addrspace(1), ptr addrspace(4) %14, align 8392 %arrayidx.i.i.i.i.i = getelementptr inbounds i32, ptr addrspace(1) %15, i64 %8393 %16 = load i32, ptr addrspace(1) %arrayidx.i.i.i.i.i, align 4394 %inc.i.i.i.i = add nsw i32 %16, 1395 store i32 %inc.i.i.i.i, ptr addrspace(1) %arrayidx.i.i.i.i.i, align 4396 %add.i.i = add i64 %storemerge.i.i, %4397 br label %for.cond.i.i398 399lexit3: ; preds = %for.cond.i.i400 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef 2, i32 noundef 2, i32 noundef 272)401 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)402 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)403 br i1 %cmpz27.i, label %TestMat.i, label %LeaderMat.i404 405TestMat.i: ; preds = %lexit3406 store i64 ptrtoint (ptr addrspace(4) addrspacecast (ptr addrspace(3) @GKernel1 to ptr addrspace(4)) to i64), ptr addrspace(3) @WGCopy.6.0, align 8407 store i64 %12, ptr addrspace(3) @WGCopy.4.0, align 8408 store i64 %11, ptr addrspace(3) @WGCopy.3.0, align 8409 store i64 2, ptr addrspace(3) @WGCopy.5.0, align 8410 br label %LeaderMat.i411 412LeaderMat.i: ; preds = %TestMat.i, %lexit3413 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)414 %17 = load i64, ptr addrspace(3) @WGCopy.5.0, align 8415 %18 = load i64, ptr addrspace(3) @WGCopy.6.0, align 8416 %19 = inttoptr i64 %18 to ptr addrspace(4)417 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)418 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef 2, i32 noundef 2, i32 noundef 272)419 %20 = getelementptr inbounds i8, ptr addrspace(4) %19, i64 24420 br label %for.cond.i.i19421 422for.cond.i.i19: ; preds = %for.body.i.i22, %LeaderMat.i423 %storemerge.i.i20 = phi i64 [ %9, %LeaderMat.i ], [ %add.i.i26, %for.body.i.i22 ]424 %cmp.i.i21 = icmp ult i64 %storemerge.i.i20, %17425 br i1 %cmp.i.i21, label %for.body.i.i22, label %lexit4426 427for.body.i.i22: ; preds = %for.cond.i.i19428 call void @llvm.assume(i1 %cmp.i.i.i.i.i.i)429 %21 = load ptr addrspace(1), ptr addrspace(4) %20, align 8430 %arrayidx.i.i.i.i.i23 = getelementptr inbounds i32, ptr addrspace(1) %21, i64 %8431 %22 = load i32, ptr addrspace(1) %arrayidx.i.i.i.i.i23, align 4432 %inc.i.i.i.i25 = add nsw i32 %22, 1433 store i32 %inc.i.i.i.i25, ptr addrspace(1) %arrayidx.i.i.i.i.i23, align 4434 %add.i.i26 = add i64 %storemerge.i.i20, %4435 br label %for.cond.i.i19436 437lexit4: ; preds = %for.cond.i.i19438 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef 2, i32 noundef 2, i32 noundef 272)439 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)440 br i1 %cmpz27.i, label %wg_leader17.i, label %wg_cf18.i441 442wg_leader17.i: ; preds = %lexit4443 %23 = load i32, ptr addrspace(3) @GCnt2, align 4444 %inc.i = add nsw i32 %23, 1445 store i32 %inc.i, ptr addrspace(3) @GCnt2, align 4446 br label %wg_cf18.i447 448wg_cf18.i: ; preds = %wg_leader17.i, %lexit4449 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)450 br label %for.cond.i451 452lexit2: ; preds = %wg_cf9.i453 call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %agg.tmp67)454 ret void455}456 457; Function Attrs: nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: write)458declare void @llvm.assume(i1 noundef)459 460; Function Attrs: convergent mustprogress norecurse nounwind461define weak_odr dso_local spir_kernel void @test1(ptr addrspace(1) noundef align 4 %_arg_dev_ptr, ptr noundef byval(%"class.sycl::_V1::range") align 8 %_arg_dev_ptr1, ptr noundef byval(%"class.sycl::_V1::range") align 8 %_arg_dev_ptr2, ptr noundef byval(%"class.sycl::_V1::range") align 8 %_arg_dev_ptr3) {462entry:463 %agg.tmp67 = alloca %"class.sycl::_V1::group", align 8464 %0 = load i64, ptr %_arg_dev_ptr1, align 8465 %1 = load i64, ptr %_arg_dev_ptr2, align 8466 %2 = load i64, ptr %_arg_dev_ptr3, align 8467 store i64 %2, ptr addrspace(3) @GKernel2, align 8468 store i64 %0, ptr addrspace(3) undef, align 8469 store i64 %1, ptr addrspace(3) undef, align 8470 %add.ptr.i = getelementptr inbounds i32, ptr addrspace(1) %_arg_dev_ptr, i64 %2471 store ptr addrspace(1) %add.ptr.i, ptr addrspace(3) undef, align 8472 %3 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalSize, align 32473 %4 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupSize, align 32474 %5 = load i64, ptr addrspace(1) @__spirv_BuiltInNumWorkgroups, align 32475 %6 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupId, align 32476 call void @llvm.lifetime.start.p0(i64 32, ptr nonnull %agg.tmp67)477 store i64 %3, ptr %agg.tmp67, align 1478 %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 8479 store i64 %4, ptr %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx, align 1480 %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 16481 store i64 %5, ptr %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx, align 1482 %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 24483 store i64 %6, ptr %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx, align 1484 %7 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationIndex, align 8485 tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)486 %cmpz15.i = icmp eq i64 %7, 0487 br i1 %cmpz15.i, label %leader.i, label %merge.i488 489leader.i: ; preds = %entry490 call void @llvm.memcpy.p3.p0.i64(ptr addrspace(3) noundef align 16 dereferenceable(32) @ArgShadow.11, ptr noundef nonnull align 8 dereferenceable(32) %agg.tmp67, i64 32, i1 false)491 br label %merge.i492 493merge.i: ; preds = %leader.i, %entry494 tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)495 call void @llvm.memcpy.p0.p3.i64(ptr noundef nonnull align 8 dereferenceable(32) %agg.tmp67, ptr addrspace(3) noundef align 16 dereferenceable(32) @ArgShadow.11, i64 32, i1 false)496 tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)497 br i1 %cmpz15.i, label %wg_leader.i, label %wg_cf.i498 499wg_leader.i: ; preds = %merge.i500 %g.ascast.i = addrspacecast ptr %agg.tmp67 to ptr addrspace(4)501 store ptr addrspace(4) %g.ascast.i, ptr addrspace(3) @GAscast3, align 8502 store i32 0, ptr addrspace(3) @GCnt3, align 4503 br label %wg_cf.i504 505wg_cf.i: ; preds = %wg_leader.i, %merge.i506 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)507 %8 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalInvocationId, align 32508 %9 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationId, align 32509 %cmp.i.i.i.i.i.i = icmp ult i64 %8, 2147483648510 br label %for.cond.i511 512for.cond.i: ; preds = %wg_cf11.i, %wg_cf.i513 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)514 br i1 %cmpz15.i, label %wg_leader4.i, label %wg_cf5.i515 516wg_leader4.i: ; preds = %for.cond.i517 %10 = load i32, ptr addrspace(3) @GCnt3, align 4518 %cmp.i = icmp slt i32 %10, 2519 store i1 %cmp.i, ptr addrspace(3) @GCmp3, align 1520 br label %wg_cf5.i521 522wg_cf5.i: ; preds = %wg_leader4.i, %for.cond.i523 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)524 %wg_val_cmp.i = load i1, ptr addrspace(3) @GCmp3, align 1525 br i1 %wg_val_cmp.i, label %for.body.i, label %lexit6526 527for.body.i: ; preds = %wg_cf5.i528 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)529 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)530 br i1 %cmpz15.i, label %TestMat.i, label %LeaderMat.i531 532TestMat.i: ; preds = %for.body.i533 store i64 ptrtoint (ptr addrspace(4) addrspacecast (ptr addrspace(3) @GKernel2 to ptr addrspace(4)) to i64), ptr addrspace(3) @WGCopy.10.0, align 8534 store i64 5, ptr addrspace(3) @WGCopy.9.0, align 8535 br label %LeaderMat.i536 537LeaderMat.i: ; preds = %TestMat.i, %for.body.i538 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)539 %11 = load i64, ptr addrspace(3) @WGCopy.9.0, align 8540 %12 = load i64, ptr addrspace(3) @WGCopy.10.0, align 8541 %13 = inttoptr i64 %12 to ptr addrspace(4)542 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)543 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef 2, i32 noundef 2, i32 noundef 272)544 %14 = getelementptr inbounds i8, ptr addrspace(4) %13, i64 24545 br label %for.cond.i.i546 547for.cond.i.i: ; preds = %for.body.i.i, %LeaderMat.i548 %storemerge.i.i = phi i64 [ %9, %LeaderMat.i ], [ %add.i.i, %for.body.i.i ]549 %cmp.i.i = icmp ult i64 %storemerge.i.i, %11550 br i1 %cmp.i.i, label %for.body.i.i, label %lexit7551 552for.body.i.i: ; preds = %for.cond.i.i553 %cmp5.not.i.i.i.i.i.i = icmp ne i64 %storemerge.i.i, %9554 %cond.i.i.i.i = zext i1 %cmp5.not.i.i.i.i.i.i to i32555 call void @llvm.assume(i1 %cmp.i.i.i.i.i.i)556 %15 = load ptr addrspace(1), ptr addrspace(4) %14, align 8557 %arrayidx.i.i.i.i.i = getelementptr inbounds i32, ptr addrspace(1) %15, i64 %8558 %16 = load i32, ptr addrspace(1) %arrayidx.i.i.i.i.i, align 4559 %add.i.i.i.i = add nsw i32 %16, %cond.i.i.i.i560 store i32 %add.i.i.i.i, ptr addrspace(1) %arrayidx.i.i.i.i.i, align 4561 %add.i.i = add i64 %storemerge.i.i, %4562 br label %for.cond.i.i563 564lexit7: ; preds = %for.cond.i.i565 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef 2, i32 noundef 2, i32 noundef 272)566 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)567 br i1 %cmpz15.i, label %wg_leader10.i, label %wg_cf11.i568 569wg_leader10.i: ; preds = %lexit7570 %17 = load i32, ptr addrspace(3) @GCnt3, align 4571 %inc.i = add nsw i32 %17, 1572 store i32 %inc.i, ptr addrspace(3) @GCnt3, align 4573 br label %wg_cf11.i574 575wg_cf11.i: ; preds = %wg_leader10.i, %lexit7576 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)577 br label %for.cond.i578 579lexit6: ; preds = %wg_cf5.i580 call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %agg.tmp67)581 ret void582}583 584; Function Attrs: convergent mustprogress norecurse nounwind585define weak_odr dso_local spir_kernel void @test2(ptr addrspace(1) noundef align 4 %_arg_dev_ptr, ptr noundef byval(%"class.sycl::_V1::range") align 8 %_arg_dev_ptr1, ptr noundef byval(%"class.sycl::_V1::range") align 8 %_arg_dev_ptr2, ptr noundef byval(%"class.sycl::_V1::range") align 8 %_arg_dev_ptr3) {586entry:587 %priv.i = alloca %"class.sycl::_V1::private_memory", align 4588 %agg.tmp67 = alloca %"class.sycl::_V1::group.15", align 8589 %0 = load i64, ptr %_arg_dev_ptr1, align 8590 %1 = load i64, ptr %_arg_dev_ptr2, align 8591 %2 = load i64, ptr %_arg_dev_ptr3, align 8592 store i64 %2, ptr addrspace(3) @GKernel3, align 8593 store i64 %0, ptr addrspace(3) undef, align 8594 store i64 %1, ptr addrspace(3) undef, align 8595 %add.ptr.i = getelementptr inbounds i32, ptr addrspace(1) %_arg_dev_ptr, i64 %2596 store ptr addrspace(1) %add.ptr.i, ptr addrspace(3) undef, align 8597 %3 = load i64, ptr addrspace(1) undef, align 8598 %4 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalSize, align 32599 %5 = load i64, ptr addrspace(1) undef, align 8600 %6 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupSize, align 32601 %7 = load i64, ptr addrspace(1) undef, align 8602 %8 = load i64, ptr addrspace(1) @__spirv_BuiltInNumWorkgroups, align 32603 %9 = load i64, ptr addrspace(1) undef, align 8604 %10 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupId, align 32605 call void @llvm.lifetime.start.p0(i64 64, ptr nonnull %agg.tmp67)606 store i64 %3, ptr %agg.tmp67, align 1607 %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 8608 store i64 %4, ptr %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx, align 1609 %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 16610 store i64 %5, ptr %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx, align 1611 %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 24612 store i64 %6, ptr %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx, align 1613 %agg.tmp6.sroa.5.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 32614 store i64 %7, ptr %agg.tmp6.sroa.5.0.agg.tmp67.sroa_idx, align 1615 %agg.tmp6.sroa.6.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 40616 store i64 %8, ptr %agg.tmp6.sroa.6.0.agg.tmp67.sroa_idx, align 1617 %agg.tmp6.sroa.7.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 48618 store i64 %9, ptr %agg.tmp6.sroa.7.0.agg.tmp67.sroa_idx, align 1619 %agg.tmp6.sroa.8.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 56620 store i64 %10, ptr %agg.tmp6.sroa.8.0.agg.tmp67.sroa_idx, align 1621 %11 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationIndex, align 8622 tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)623 %cmpz32.i = icmp eq i64 %11, 0624 br i1 %cmpz32.i, label %leader.i, label %merge.i625 626leader.i: ; preds = %entry627 call void @llvm.memcpy.p3.p0.i64(ptr addrspace(3) noundef align 16 dereferenceable(64) @ArgShadow.17, ptr noundef nonnull align 8 dereferenceable(64) %agg.tmp67, i64 64, i1 false)628 br label %merge.i629 630merge.i: ; preds = %leader.i, %entry631 tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)632 call void @llvm.memcpy.p0.p3.i64(ptr noundef nonnull align 8 dereferenceable(64) %agg.tmp67, ptr addrspace(3) noundef align 16 dereferenceable(64) @ArgShadow.17, i64 64, i1 false)633 %priv.ascast.i = addrspacecast ptr %priv.i to ptr addrspace(4)634 tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)635 br i1 %cmpz32.i, label %wg_leader.i, label %wg_cf.i636 637wg_leader.i: ; preds = %merge.i638 %g.ascast.i = addrspacecast ptr %agg.tmp67 to ptr addrspace(4)639 store ptr addrspace(4) %g.ascast.i, ptr addrspace(3) @GAsCast4, align 8640 call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %priv.i)641 store i32 0, ptr addrspace(3) @GCnt4, align 4642 br label %wg_cf.i643 644wg_cf.i: ; preds = %wg_leader.i, %merge.i645 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)646 %12 = load i64, ptr addrspace(1) undef, align 8647 %13 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalInvocationId, align 32648 %14 = load i64, ptr addrspace(1) undef, align 8649 %15 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationId, align 32650 %mul.i.i.i.i.i.i = mul i64 %12, %4651 %add.i.i.i.i.i.i = add i64 %mul.i.i.i.i.i.i, %13652 %cmp.i.i.i.i.i.i = icmp ult i64 %add.i.i.i.i.i.i, 2147483648653 %conv.i.i.i.i.i = trunc i64 %add.i.i.i.i.i.i to i32654 %y.i.i.i.i.i = getelementptr inbounds i8, ptr %priv.i, i64 4655 br label %for.cond.i656 657for.cond.i: ; preds = %wg_cf20.i, %wg_cf.i658 %agg.tmp6.i.sroa.9.0 = phi ptr addrspace(4) [ undef, %wg_cf.i ], [ %agg.tmp6.i.sroa.9.0.copyload40, %wg_cf20.i ]659 %agg.tmp5.i.sroa.0.0 = phi i64 [ undef, %wg_cf.i ], [ %agg.tmp5.i.sroa.0.0.copyload44, %wg_cf20.i ]660 %agg.tmp5.i.sroa.8.0 = phi i64 [ undef, %wg_cf.i ], [ %agg.tmp5.i.sroa.8.0.copyload48, %wg_cf20.i ]661 %agg.tmp2.i.sroa.0.0 = phi ptr addrspace(4) [ undef, %wg_cf.i ], [ %agg.tmp2.i.sroa.0.0.copyload52, %wg_cf20.i ]662 %agg.tmp2.i.sroa.8.0 = phi ptr addrspace(4) [ undef, %wg_cf.i ], [ %agg.tmp2.i.sroa.8.0.copyload56, %wg_cf20.i ]663 %agg.tmp.i.sroa.0.0 = phi i64 [ undef, %wg_cf.i ], [ %agg.tmp.i.sroa.0.0.copyload60, %wg_cf20.i ]664 %agg.tmp.i.sroa.8.0 = phi i64 [ undef, %wg_cf.i ], [ %agg.tmp.i.sroa.8.0.copyload64, %wg_cf20.i ]665 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)666 br i1 %cmpz32.i, label %wg_leader10.i, label %wg_cf11.i667 668wg_leader10.i: ; preds = %for.cond.i669 %16 = load i32, ptr addrspace(3) @GCnt4, align 4670 %cmp.i = icmp slt i32 %16, 2671 store i1 %cmp.i, ptr addrspace(3) @GCmp4, align 1672 br label %wg_cf11.i673 674wg_cf11.i: ; preds = %wg_leader10.i, %for.cond.i675 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)676 %wg_val_cmp.i = load i1, ptr addrspace(3) @GCmp4, align 1677 br i1 %wg_val_cmp.i, label %for.body.i, label %for.end.i678 679for.body.i: ; preds = %wg_cf11.i680 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)681 br i1 %cmpz32.i, label %wg_leader13.i, label %wg_cf14.i682 683wg_leader13.i: ; preds = %for.body.i684 br label %wg_cf14.i685 686wg_cf14.i: ; preds = %wg_leader13.i, %for.body.i687 %agg.tmp2.i.sroa.0.1 = phi ptr addrspace(4) [ addrspacecast (ptr addrspace(3) @GKernel3 to ptr addrspace(4)), %wg_leader13.i ], [ %agg.tmp2.i.sroa.0.0, %for.body.i ]688 %agg.tmp2.i.sroa.8.1 = phi ptr addrspace(4) [ %priv.ascast.i, %wg_leader13.i ], [ %agg.tmp2.i.sroa.8.0, %for.body.i ]689 %agg.tmp.i.sroa.0.1 = phi i64 [ 7, %wg_leader13.i ], [ %agg.tmp.i.sroa.0.0, %for.body.i ]690 %agg.tmp.i.sroa.8.1 = phi i64 [ 3, %wg_leader13.i ], [ %agg.tmp.i.sroa.8.0, %for.body.i ]691 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)692 br i1 %cmpz32.i, label %TestMat30.i, label %LeaderMat27.i693 694TestMat30.i: ; preds = %wg_cf14.i695 store i64 %agg.tmp.i.sroa.0.1, ptr addrspace(3) @WGCopy.13.0, align 8696 store i64 %agg.tmp.i.sroa.8.1, ptr addrspace(3) @WGCopy.13.1, align 8697 store ptr addrspace(4) %agg.tmp2.i.sroa.0.1, ptr addrspace(3) @WGCopy.14.0, align 8698 store ptr addrspace(4) %agg.tmp2.i.sroa.8.1, ptr addrspace(3) @WGCopy.14.1, align 8699 store i64 %agg.tmp5.i.sroa.0.0, ptr addrspace(3) @WGCopy.15.0, align 8700 store i64 %agg.tmp5.i.sroa.8.0, ptr addrspace(3) @WGCopy.15.1, align 8701 store ptr addrspace(4) %priv.ascast.i, ptr addrspace(3) @WGCopy.16.0, align 8702 store ptr addrspace(4) %agg.tmp6.i.sroa.9.0, ptr addrspace(3) @WGCopy.16.1, align 8703 br label %LeaderMat27.i704 705LeaderMat27.i: ; preds = %TestMat30.i, %wg_cf14.i706 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)707 %agg.tmp6.i.sroa.0.0.copyload = load ptr addrspace(4), ptr addrspace(3) @WGCopy.16.0, align 8708 %agg.tmp6.i.sroa.9.0.copyload = load ptr addrspace(4), ptr addrspace(3) @WGCopy.16.1, align 8709 %agg.tmp5.i.sroa.0.0.copyload = load i64, ptr addrspace(3) @WGCopy.15.0, align 8710 %agg.tmp5.i.sroa.8.0.copyload = load i64, ptr addrspace(3) @WGCopy.15.1, align 8711 %agg.tmp2.i.sroa.0.0.copyload = load ptr addrspace(4), ptr addrspace(3) @WGCopy.14.0, align 8712 %agg.tmp.i.sroa.0.0.copyload = load i64, ptr addrspace(3) @WGCopy.13.0, align 8713 %agg.tmp.i.sroa.8.0.copyload = load i64, ptr addrspace(3) @WGCopy.13.1, align 8714 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)715 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef 2, i32 noundef 2, i32 noundef 272)716 %17 = getelementptr inbounds i8, ptr addrspace(4) %agg.tmp2.i.sroa.0.0.copyload, i64 24717 br label %for.cond.i.i718 719for.cond.i.i: ; preds = %lexit10, %LeaderMat27.i720 %storemerge.i.i = phi i64 [ %14, %LeaderMat27.i ], [ %add.i.i, %lexit10 ]721 %cmp.i.i = icmp ult i64 %storemerge.i.i, %agg.tmp.i.sroa.0.0.copyload722 br i1 %cmp.i.i, label %for.cond.i.i.i, label %lexit11723 724for.cond.i.i.i: ; preds = %for.body.i.i.i, %for.cond.i.i725 %storemerge.i.i.i = phi i64 [ %add.i.i.i, %for.body.i.i.i ], [ %15, %for.cond.i.i ]726 %cmp.i.i.i = icmp ult i64 %storemerge.i.i.i, %agg.tmp.i.sroa.8.0.copyload727 br i1 %cmp.i.i.i, label %for.body.i.i.i, label %lexit10728 729for.body.i.i.i: ; preds = %for.cond.i.i.i730 call void @llvm.assume(i1 %cmp.i.i.i.i.i.i)731 %18 = load ptr addrspace(1), ptr addrspace(4) %17, align 8732 %arrayidx.i.i.i.i.i.i = getelementptr inbounds i32, ptr addrspace(1) %18, i64 %add.i.i.i.i.i.i733 %19 = load i32, ptr addrspace(1) %arrayidx.i.i.i.i.i.i, align 4734 %inc.i.i.i.i.i = add nsw i32 %19, 1735 store i32 %inc.i.i.i.i.i, ptr addrspace(1) %arrayidx.i.i.i.i.i.i, align 4736 store i32 %conv.i.i.i.i.i, ptr %priv.i, align 4737 store i32 5, ptr %y.i.i.i.i.i, align 4738 %add.i.i.i = add i64 %storemerge.i.i.i, %6739 br label %for.cond.i.i.i740 741lexit10: ; preds = %for.cond.i.i.i742 %add.i.i = add i64 %storemerge.i.i, %5743 br label %for.cond.i.i744 745lexit11: ; preds = %for.cond.i.i746 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef 2, i32 noundef 2, i32 noundef 272)747 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)748 br i1 %cmpz32.i, label %wg_leader16.i, label %wg_cf17.i749 750wg_leader16.i: ; preds = %lexit11751 br label %wg_cf17.i752 753wg_cf17.i: ; preds = %wg_leader16.i, %lexit11754 %agg.tmp6.i.sroa.0.1 = phi ptr addrspace(4) [ %priv.ascast.i, %wg_leader16.i ], [ %agg.tmp6.i.sroa.0.0.copyload, %lexit11 ]755 %agg.tmp6.i.sroa.9.1 = phi ptr addrspace(4) [ addrspacecast (ptr addrspace(3) @GKernel3 to ptr addrspace(4)), %wg_leader16.i ], [ %agg.tmp6.i.sroa.9.0.copyload, %lexit11 ]756 %agg.tmp5.i.sroa.0.1 = phi i64 [ 7, %wg_leader16.i ], [ %agg.tmp5.i.sroa.0.0.copyload, %lexit11 ]757 %agg.tmp5.i.sroa.8.1 = phi i64 [ 3, %wg_leader16.i ], [ %agg.tmp5.i.sroa.8.0.copyload, %lexit11 ]758 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)759 br i1 %cmpz32.i, label %TestMat.i, label %LeaderMat.i760 761TestMat.i: ; preds = %wg_cf17.i762 store i64 %agg.tmp.i.sroa.0.0.copyload, ptr addrspace(3) @WGCopy.13.0, align 8763 store i64 %agg.tmp.i.sroa.8.0.copyload, ptr addrspace(3) @WGCopy.13.1, align 8764 store ptr addrspace(4) %agg.tmp2.i.sroa.0.0.copyload, ptr addrspace(3) @WGCopy.14.0, align 8765 store ptr addrspace(4) %priv.ascast.i, ptr addrspace(3) @WGCopy.14.1, align 8766 store i64 %agg.tmp5.i.sroa.0.1, ptr addrspace(3) @WGCopy.15.0, align 8767 store i64 %agg.tmp5.i.sroa.8.1, ptr addrspace(3) @WGCopy.15.1, align 8768 store ptr addrspace(4) %agg.tmp6.i.sroa.0.1, ptr addrspace(3) @WGCopy.16.0, align 8769 store ptr addrspace(4) %agg.tmp6.i.sroa.9.1, ptr addrspace(3) @WGCopy.16.1, align 8770 br label %LeaderMat.i771 772LeaderMat.i: ; preds = %TestMat.i, %wg_cf17.i773 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)774 %agg.tmp6.i.sroa.9.0.copyload40 = load ptr addrspace(4), ptr addrspace(3) @WGCopy.16.1, align 8775 %agg.tmp5.i.sroa.0.0.copyload44 = load i64, ptr addrspace(3) @WGCopy.15.0, align 8776 %agg.tmp5.i.sroa.8.0.copyload48 = load i64, ptr addrspace(3) @WGCopy.15.1, align 8777 %agg.tmp2.i.sroa.0.0.copyload52 = load ptr addrspace(4), ptr addrspace(3) @WGCopy.14.0, align 8778 %agg.tmp2.i.sroa.8.0.copyload56 = load ptr addrspace(4), ptr addrspace(3) @WGCopy.14.1, align 8779 %agg.tmp.i.sroa.0.0.copyload60 = load i64, ptr addrspace(3) @WGCopy.13.0, align 8780 %agg.tmp.i.sroa.8.0.copyload64 = load i64, ptr addrspace(3) @WGCopy.13.1, align 8781 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)782 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef 2, i32 noundef 2, i32 noundef 272)783 %20 = getelementptr inbounds i8, ptr addrspace(4) %agg.tmp6.i.sroa.9.0.copyload40, i64 24784 br label %for.cond.i.i25785 786for.cond.i.i25: ; preds = %lexit12, %LeaderMat.i787 %storemerge.i.i26 = phi i64 [ %14, %LeaderMat.i ], [ %add.i.i31, %lexit12 ]788 %cmp.i.i27 = icmp ult i64 %storemerge.i.i26, %agg.tmp5.i.sroa.0.0.copyload44789 br i1 %cmp.i.i27, label %for.cond.i.i.i28, label %lexit13790 791for.cond.i.i.i28: ; preds = %for.body.i.i.i32, %for.cond.i.i25792 %storemerge.i.i.i29 = phi i64 [ %add.i.i.i35, %for.body.i.i.i32 ], [ %15, %for.cond.i.i25 ]793 %cmp.i.i.i30 = icmp ult i64 %storemerge.i.i.i29, %agg.tmp5.i.sroa.8.0.copyload48794 br i1 %cmp.i.i.i30, label %for.body.i.i.i32, label %lexit12795 796for.body.i.i.i32: ; preds = %for.cond.i.i.i28797 %21 = load i32, ptr %priv.i, align 4798 %22 = load i32, ptr %y.i.i.i.i.i, align 4799 %add.i.i.i.i.i = add nsw i32 %21, %22800 call void @llvm.assume(i1 %cmp.i.i.i.i.i.i)801 %23 = load ptr addrspace(1), ptr addrspace(4) %20, align 8802 %arrayidx.i.i.i.i.i.i33 = getelementptr inbounds i32, ptr addrspace(1) %23, i64 %add.i.i.i.i.i.i803 %24 = load i32, ptr addrspace(1) %arrayidx.i.i.i.i.i.i33, align 4804 %add4.i.i.i.i.i = add nsw i32 %24, %add.i.i.i.i.i805 store i32 %add4.i.i.i.i.i, ptr addrspace(1) %arrayidx.i.i.i.i.i.i33, align 4806 %add.i.i.i35 = add i64 %storemerge.i.i.i29, %6807 br label %for.cond.i.i.i28808 809lexit12: ; preds = %for.cond.i.i.i28810 %add.i.i31 = add i64 %storemerge.i.i26, %5811 br label %for.cond.i.i25812 813lexit13: ; preds = %for.cond.i.i25814 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef 2, i32 noundef 2, i32 noundef 272)815 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)816 br i1 %cmpz32.i, label %wg_leader19.i, label %wg_cf20.i817 818wg_leader19.i: ; preds = %lexit13819 %25 = load i32, ptr addrspace(3) @GCnt4, align 4820 %inc.i = add nsw i32 %25, 1821 store i32 %inc.i, ptr addrspace(3) @GCnt4, align 4822 br label %wg_cf20.i823 824wg_cf20.i: ; preds = %wg_leader19.i, %lexit13825 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)826 br label %for.cond.i827 828for.end.i: ; preds = %wg_cf11.i829 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)830 br i1 %cmpz32.i, label %wg_leader22.i, label %lexit14831 832wg_leader22.i: ; preds = %for.end.i833 call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %priv.i)834 br label %lexit14835 836lexit14: ; preds = %wg_leader22.i, %for.end.i837 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)838 call void @llvm.lifetime.end.p0(i64 64, ptr nonnull %agg.tmp67)839 ret void840}841 842; Function Attrs: convergent mustprogress norecurse nounwind843define weak_odr dso_local spir_kernel void @test3(ptr addrspace(1) noundef align 4 %_arg_dev_ptr, ptr noundef byval(%"class.sycl::_V1::range") align 8 %_arg_dev_ptr1, ptr noundef byval(%"class.sycl::_V1::range") align 8 %_arg_dev_ptr2, ptr noundef byval(%"class.sycl::_V1::range") align 8 %_arg_dev_ptr3) {844entry:845 %agg.tmp67 = alloca %"class.sycl::_V1::group", align 8846 %0 = load i64, ptr %_arg_dev_ptr1, align 8847 %1 = load i64, ptr %_arg_dev_ptr2, align 8848 %2 = load i64, ptr %_arg_dev_ptr3, align 8849 store i64 %2, ptr addrspace(3) @GKernel4, align 8850 store i64 %0, ptr addrspace(3) undef, align 8851 store i64 %1, ptr addrspace(3) undef, align 8852 %add.ptr.i = getelementptr inbounds i32, ptr addrspace(1) %_arg_dev_ptr, i64 %2853 store ptr addrspace(1) %add.ptr.i, ptr addrspace(3) undef, align 8854 %3 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalSize, align 32855 %4 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupSize, align 32856 %5 = load i64, ptr addrspace(1) @__spirv_BuiltInNumWorkgroups, align 32857 %6 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupId, align 32858 call void @llvm.lifetime.start.p0(i64 32, ptr nonnull %agg.tmp67)859 store i64 %3, ptr %agg.tmp67, align 1860 %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 8861 store i64 %4, ptr %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx, align 1862 %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 16863 store i64 %5, ptr %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx, align 1864 %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 24865 store i64 %6, ptr %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx, align 1866 %7 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationIndex, align 8867 tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)868 %cmpz16.i = icmp eq i64 %7, 0869 br i1 %cmpz16.i, label %leader.i, label %merge.i870 871leader.i: ; preds = %entry872 call void @llvm.memcpy.p3.p0.i64(ptr addrspace(3) noundef align 16 dereferenceable(32) @ArgShadow.21, ptr noundef nonnull align 8 dereferenceable(32) %agg.tmp67, i64 32, i1 false)873 br label %merge.i874 875merge.i: ; preds = %leader.i, %entry876 tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)877 call void @llvm.memcpy.p0.p3.i64(ptr noundef nonnull align 8 dereferenceable(32) %agg.tmp67, ptr addrspace(3) noundef align 16 dereferenceable(32) @ArgShadow.21, i64 32, i1 false)878 tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)879 br i1 %cmpz16.i, label %wg_leader.i, label %wg_cf.i880 881wg_leader.i: ; preds = %merge.i882 %g.ascast.i = addrspacecast ptr %agg.tmp67 to ptr addrspace(4)883 store ptr addrspace(4) %g.ascast.i, ptr addrspace(3) @GAsCast5, align 8884 store i32 0, ptr addrspace(3) @GCnt5, align 4885 br label %wg_cf.i886 887wg_cf.i: ; preds = %wg_leader.i, %merge.i888 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)889 %wg_val_g.ascast.i = load ptr addrspace(4), ptr addrspace(3) @GAsCast5, align 8890 %8 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationId, align 32891 %9 = trunc i64 %4 to i32892 br label %for.cond.i893 894for.cond.i: ; preds = %wg_cf12.i, %wg_cf.i895 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)896 br i1 %cmpz16.i, label %wg_leader5.i, label %wg_cf6.i897 898wg_leader5.i: ; preds = %for.cond.i899 %10 = load i32, ptr addrspace(3) @GCnt5, align 4900 %cmp.i = icmp slt i32 %10, 2901 store i1 %cmp.i, ptr addrspace(3) @GCmp5, align 1902 br label %wg_cf6.i903 904wg_cf6.i: ; preds = %wg_leader5.i, %for.cond.i905 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)906 %wg_val_cmp.i = load i1, ptr addrspace(3) @GCmp5, align 1907 br i1 %wg_val_cmp.i, label %for.body.i, label %lexit20908 909for.body.i: ; preds = %wg_cf6.i910 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)911 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)912 br i1 %cmpz16.i, label %TestMat.i, label %LeaderMat.i913 914TestMat.i: ; preds = %for.body.i915 store ptr addrspace(4) %wg_val_g.ascast.i, ptr addrspace(3) @WGCopy.20.0, align 8916 store ptr addrspace(4) addrspacecast (ptr addrspace(3) @GKernel4 to ptr addrspace(4)), ptr addrspace(3) @WGCopy.20.1, align 8917 store i64 5, ptr addrspace(3) @WGCopy.19.0, align 8918 br label %LeaderMat.i919 920LeaderMat.i: ; preds = %TestMat.i, %for.body.i921 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)922 %11 = load i64, ptr addrspace(3) @WGCopy.19.0, align 8923 %agg.tmp2.i.sroa.0.0.copyload = load ptr addrspace(4), ptr addrspace(3) @WGCopy.20.0, align 8924 %agg.tmp2.i.sroa.6.0.copyload = load ptr addrspace(4), ptr addrspace(3) @WGCopy.20.1, align 8925 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)926 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef 2, i32 noundef 2, i32 noundef 272)927 %index.i.i.i.i.i = getelementptr inbounds i8, ptr addrspace(4) %agg.tmp2.i.sroa.0.0.copyload, i64 24928 %12 = getelementptr inbounds i8, ptr addrspace(4) %agg.tmp2.i.sroa.6.0.copyload, i64 24929 %13 = trunc i64 %11 to i32930 br label %for.cond.i.i931 932for.cond.i.i: ; preds = %for.body.i.i, %LeaderMat.i933 %storemerge.i.i = phi i64 [ %8, %LeaderMat.i ], [ %add.i.i, %for.body.i.i ]934 %cmp.i.i = icmp ult i64 %storemerge.i.i, %11935 br i1 %cmp.i.i, label %for.body.i.i, label %lexit21936 937for.body.i.i: ; preds = %for.cond.i.i938 %14 = load i64, ptr addrspace(4) %index.i.i.i.i.i, align 8939 %mul.i.i.i.i = mul i64 %14, 10940 %mul3.i.i.i.i = shl i64 %storemerge.i.i, 1941 %add.i.i.i.i = add i64 %mul.i.i.i.i, %mul3.i.i.i.i942 %15 = load ptr addrspace(1), ptr addrspace(4) %12, align 8943 %arrayidx.i.i.i.i.i = getelementptr inbounds i32, ptr addrspace(1) %15, i64 %add.i.i.i.i944 %16 = load i32, ptr addrspace(1) %arrayidx.i.i.i.i.i, align 4945 %conv9.i.i.i.i = add i32 %16, %13946 store i32 %conv9.i.i.i.i, ptr addrspace(1) %arrayidx.i.i.i.i.i, align 4947 %add14.i.i.i.i = or disjoint i64 %add.i.i.i.i, 1948 %17 = load ptr addrspace(1), ptr addrspace(4) %12, align 8949 %arrayidx.i25.i.i.i.i = getelementptr inbounds i32, ptr addrspace(1) %17, i64 %add14.i.i.i.i950 %18 = load i32, ptr addrspace(1) %arrayidx.i25.i.i.i.i, align 4951 %conv18.i.i.i.i = add i32 %18, %9952 store i32 %conv18.i.i.i.i, ptr addrspace(1) %arrayidx.i25.i.i.i.i, align 4953 %add.i.i = add i64 %storemerge.i.i, %4954 br label %for.cond.i.i955 956lexit21: ; preds = %for.cond.i.i957 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef 2, i32 noundef 2, i32 noundef 272)958 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)959 br i1 %cmpz16.i, label %wg_leader11.i, label %wg_cf12.i960 961wg_leader11.i: ; preds = %lexit21962 %19 = load i32, ptr addrspace(3) @GCnt5, align 4963 %inc.i = add nsw i32 %19, 1964 store i32 %inc.i, ptr addrspace(3) @GCnt5, align 4965 br label %wg_cf12.i966 967wg_cf12.i: ; preds = %wg_leader11.i, %lexit21968 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)969 br label %for.cond.i970 971lexit20: ; preds = %wg_cf6.i972 call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %agg.tmp67)973 ret void974}975