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 8 22@WI.0 = internal unnamed_addr addrspace(3) global i64 undef, align 8 23@WI.1 = internal unnamed_addr addrspace(3) global i64 undef, align 8 24@WI.2 = internal unnamed_addr addrspace(3) global i64 undef, align 8 25@WI.3 = internal unnamed_addr addrspace(3) global i64 undef, align 8 26@WI.4 = internal unnamed_addr addrspace(3) global i32 undef, align 8 27@WI.6 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 8 28@GCnt = internal unnamed_addr addrspace(3) global i32 undef, align 4 29@__spirv_BuiltInNumWorkgroups = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32 30@GKernel1 = internal addrspace(3) global %class.anon.2 undef, align 8 31@GCnt2 = internal unnamed_addr addrspace(3) global i32 undef, align 4 32@GKernel2 = internal addrspace(3) global %class.anon.2 undef, align 8 33@GCnt3 = internal unnamed_addr addrspace(3) global i32 undef, align 4 34@GKernel3 = internal addrspace(3) global %class.anon.2 undef, align 8 35@GCnt4 = internal unnamed_addr addrspace(3) global i32 undef, align 4 36@GKernel4 = internal addrspace(3) global %class.anon.2 undef, align 8 37@GCnt5 = internal unnamed_addr addrspace(3) global i32 undef, align 4 38@__spirv_BuiltInLocalInvocationIndex = external local_unnamed_addr addrspace(1) constant i64, align 8 39@GThis = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 8 40@GAsCast = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 8 41@GCmp = internal unnamed_addr addrspace(3) global i1 undef, align 1 42@WGCopy = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 8 43@WGCopy.1.0 = internal unnamed_addr addrspace(3) global i64 undef, align 16 44@WGCopy.1.1 = internal unnamed_addr addrspace(3) global i64 undef, align 16 45@WGCopy.1.2 = internal unnamed_addr addrspace(3) global i64 undef, align 16 46@WGCopy.1.3 = internal unnamed_addr addrspace(3) global i64 undef, align 16 47@WGCopy.1.4 = internal unnamed_addr addrspace(3) global i32 undef, align 16 48@WGCopy.1.5 = internal unnamed_addr addrspace(3) global i32 undef, align 16 49@WGCopy.1.6 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 16 50@ArgShadow = internal unnamed_addr addrspace(3) global %"class.sycl::_V1::group" undef, align 16 51@GAsCast2 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 8 52@GCmp2 = internal unnamed_addr addrspace(3) global i1 undef, align 1 53@WGCopy.3.0 = internal unnamed_addr addrspace(3) global i64 undef, align 8 54@WGCopy.4.0 = internal unnamed_addr addrspace(3) global i64 undef, align 8 55@WGCopy.5.0 = internal unnamed_addr addrspace(3) global i64 undef, align 8 56@WGCopy.6.0 = internal unnamed_addr addrspace(3) global i64 undef, align 8 57@ArgShadow.7 = internal unnamed_addr addrspace(3) global %"class.sycl::_V1::group" undef, align 16 58@GAscast3 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 8 59@GCmp3 = internal unnamed_addr addrspace(3) global i1 undef, align 1 60@WGCopy.9.0 = internal unnamed_addr addrspace(3) global i64 undef, align 8 61@WGCopy.10.0 = internal unnamed_addr addrspace(3) global i64 undef, align 8 62@ArgShadow.11 = internal unnamed_addr addrspace(3) global %"class.sycl::_V1::group" undef, align 16 63@GAsCast4 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 8 64@GCmp4 = internal unnamed_addr addrspace(3) global i1 undef, align 1 65@WGCopy.13.0 = internal unnamed_addr addrspace(3) global i64 undef, align 8 66@WGCopy.13.1 = internal unnamed_addr addrspace(3) global i64 undef, align 8 67@WGCopy.14.0 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 8 68@WGCopy.14.1 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 8 69@WGCopy.15.0 = internal unnamed_addr addrspace(3) global i64 undef, align 8 70@WGCopy.15.1 = internal unnamed_addr addrspace(3) global i64 undef, align 8 71@WGCopy.16.0 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 8 72@WGCopy.16.1 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 8 73@ArgShadow.17 = internal unnamed_addr addrspace(3) global %"class.sycl::_V1::group.15" undef, align 16 74@GAsCast5 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 8 75@GCmp5 = internal unnamed_addr addrspace(3) global i1 undef, align 1 76@WGCopy.19.0 = internal unnamed_addr addrspace(3) global i64 undef, align 8 77@WGCopy.20.0 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 8 78@WGCopy.20.1 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 8 79@ArgShadow.21 = internal unnamed_addr addrspace(3) global %"class.sycl::_V1::group" undef, align 16 80@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32 81@__spirv_BuiltInGlobalSize = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32 82@__spirv_BuiltInLocalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32 83@__spirv_BuiltInWorkgroupId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32 84@__spirv_BuiltInWorkgroupSize = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32 85 86; Function Attrs: convergent mustprogress norecurse nounwind 87define 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 8 90 store i64 %_arg_wg_chunk, ptr addrspace(3) @GFunctor, align 8 91 store i64 %_arg_range_length, ptr addrspace(3) undef, align 8 92 store i32 %_arg_n_iter, ptr addrspace(3) undef, align 8 93 store i32 %_arg_addend, ptr addrspace(3) undef, align 4 94 %0 = load i64, ptr %_arg_dev_ptr1, align 8 95 %1 = load i64, ptr %_arg_dev_ptr2, align 8 96 %2 = load i64, ptr %_arg_dev_ptr3, align 8 97 store i64 %2, ptr addrspace(3) undef, align 8 98 store i64 %0, ptr addrspace(3) undef, align 8 99 store i64 %1, ptr addrspace(3) undef, align 8 100 %add.ptr.i = getelementptr inbounds i32, ptr addrspace(1) %_arg_dev_ptr, i64 %2 101 store ptr addrspace(1) %add.ptr.i, ptr addrspace(3) undef, align 8 102 %3 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalSize, align 32 103 %4 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupSize, align 32 104 %5 = load i64, ptr addrspace(1) @__spirv_BuiltInNumWorkgroups, align 32 105 %6 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupId, align 32 106 call void @llvm.lifetime.start.p0(i64 32, ptr nonnull %agg.tmp67) 107 store i64 %3, ptr %agg.tmp67, align 1 108 %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 8 109 store i64 %4, ptr %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx, align 1 110 %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 16 111 store i64 %5, ptr %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx, align 1 112 %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 24 113 store i64 %6, ptr %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx, align 1 114 %7 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationIndex, align 8 115 tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272) 116 %cmpz15.i = icmp eq i64 %7, 0 117 br i1 %cmpz15.i, label %leader.i, label %merge.i 118 119leader.i: ; preds = %entry 120 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.i 122 123merge.i: ; preds = %leader.i, %entry 124 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.i 128 129wg_leader.i: ; preds = %merge.i 130 %g.ascast.i = addrspacecast ptr %agg.tmp67 to ptr addrspace(4) 131 store ptr addrspace(4) %g.ascast.i, ptr addrspace(3) @GAsCast, align 8 132 store ptr addrspace(4) addrspacecast (ptr addrspace(3) @GFunctor to ptr addrspace(4)), ptr addrspace(3) @GThis, align 8 133 %8 = load i32, ptr addrspace(3) undef, align 4 134 %9 = load i64, ptr addrspace(3) @GFunctor, align 8 135 %index.i = getelementptr inbounds i8, ptr %agg.tmp67, i64 24 136 %10 = load i64, ptr %index.i, align 8 137 %mul.i = mul i64 %9, %10 138 %localRange.i = getelementptr inbounds i8, ptr %agg.tmp67, i64 8 139 %11 = load i64, ptr %localRange.i, align 8 140 %12 = load i64, ptr addrspace(3) undef, align 8 141 store i64 %9, ptr addrspace(3) @WI.0, align 8 142 store i64 %11, ptr addrspace(3) @WI.1, align 8 143 store i64 %mul.i, ptr addrspace(3) @WI.2, align 8 144 store i64 %12, ptr addrspace(3) @WI.3, align 8 145 store i32 %8, ptr addrspace(3) @WI.4, align 8 146 store ptr addrspace(4) undef, ptr addrspace(3) @WI.6, align 8 147 store i32 0, ptr addrspace(3) @GCnt, align 4 148 br label %wg_cf.i 149 150wg_cf.i: ; preds = %wg_leader.i, %merge.i 151 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 8 153 %n_iter.i = getelementptr inbounds i8, ptr addrspace(4) %wg_val_this1.i, i64 16 154 %13 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationId, align 32 155 br label %for.cond.i 156 157for.cond.i: ; preds = %wg_cf11.i, %wg_cf.i 158 %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.i 168 169wg_leader4.i: ; preds = %for.cond.i 170 %14 = load i32, ptr addrspace(3) @GCnt, align 4 171 %15 = load i32, ptr addrspace(4) %n_iter.i, align 8 172 %cmp.i = icmp slt i32 %14, %15 173 store i1 %cmp.i, ptr addrspace(3) @GCmp, align 1 174 br label %wg_cf5.i 175 176wg_cf5.i: ; preds = %wg_leader4.i, %for.cond.i 177 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 1 179 br i1 %wg_val_cmp.i, label %for.body.i, label %_ZNK11PFWGFunctorclEN4sycl3_V15groupILi1EEE.exit 180 181for.body.i: ; preds = %wg_cf5.i 182 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.i 184 185wg_leader7.i: ; preds = %for.body.i 186 %agg.tmp.i.sroa.0.0.copyload = load i64, ptr addrspace(3) @WI.0, align 8 187 %agg.tmp.i.sroa.6.0.copyload = load i64, ptr addrspace(3) @WI.1, align 8 188 %agg.tmp.i.sroa.7.0.copyload = load i64, ptr addrspace(3) @WI.2, align 8 189 %agg.tmp.i.sroa.8.0.copyload = load i64, ptr addrspace(3) @WI.3, align 8 190 %agg.tmp.i.sroa.9.0.copyload = load i32, ptr addrspace(3) @WI.4, align 8 191 %agg.tmp.i.sroa.11.0.copyload = load ptr addrspace(4), ptr addrspace(3) @WI.6, align 8 192 br label %wg_cf8.i 193 194wg_cf8.i: ; preds = %wg_leader7.i, %for.body.i 195 %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.i 203 204TestMat.i: ; preds = %wg_cf8.i 205 store i64 %agg.tmp.i.sroa.0.1, ptr addrspace(3) @WGCopy.1.0, align 16 206 store i64 %agg.tmp.i.sroa.6.1, ptr addrspace(3) @WGCopy.1.1, align 16 207 store i64 %agg.tmp.i.sroa.7.1, ptr addrspace(3) @WGCopy.1.2, align 16 208 store i64 %agg.tmp.i.sroa.8.1, ptr addrspace(3) @WGCopy.1.3, align 16 209 store i32 %agg.tmp.i.sroa.9.1, ptr addrspace(3) @WGCopy.1.4, align 16 210 store i32 %agg.tmp.i.sroa.10.0, ptr addrspace(3) @WGCopy.1.5, align 16 211 store ptr addrspace(4) %agg.tmp.i.sroa.11.1, ptr addrspace(3) @WGCopy.1.6, align 16 212 store ptr addrspace(4) %this.addr.0.i, ptr addrspace(3) @WGCopy, align 8 213 br label %LeaderMat.i 214 215LeaderMat.i: ; preds = %TestMat.i, %wg_cf8.i 216 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 8 218 %agg.tmp.i.sroa.0.0.copyload13 = load i64, ptr addrspace(3) @WGCopy.1.0, align 16 219 %agg.tmp.i.sroa.6.0.copyload15 = load i64, ptr addrspace(3) @WGCopy.1.1, align 16 220 %agg.tmp.i.sroa.7.0.copyload17 = load i64, ptr addrspace(3) @WGCopy.1.2, align 16 221 %agg.tmp.i.sroa.8.0.copyload19 = load i64, ptr addrspace(3) @WGCopy.1.3, align 16 222 %agg.tmp.i.sroa.9.0.copyload21 = load i32, ptr addrspace(3) @WGCopy.1.4, align 16 223 %agg.tmp.i.sroa.10.0.copyload23 = load i32, ptr addrspace(3) @WGCopy.1.5, align 16 224 %agg.tmp.i.sroa.11.0.copyload25 = load ptr addrspace(4), ptr addrspace(3) @WGCopy.1.6, align 16 225 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.copyload13 228 br i1 %cmp.not.i.i, label %if.end.i.i, label %lexit1 229 230if.end.i.i: ; preds = %LeaderMat.i 231 %add.i.i = add i64 %agg.tmp.i.sroa.0.0.copyload13, %agg.tmp.i.sroa.6.0.copyload15 232 %sub.i.i = add i64 %add.i.i, -1 233 %div.i.i = udiv i64 %sub.i.i, %agg.tmp.i.sroa.6.0.copyload15 234 %mul.i.i = mul i64 %13, %div.i.i 235 %add4.i.i = add i64 %agg.tmp.i.sroa.7.0.copyload17, %mul.i.i 236 %add6.i.i = add i64 %add4.i.i, %div.i.i 237 %.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 24 239 br label %for.cond.i.i 240 241for.cond.i.i: ; preds = %for.body.i.i, %if.end.i.i 242 %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.i 244 br i1 %cmp8.i.i, label %for.body.i.i, label %lexit1 245 246for.body.i.i: ; preds = %for.cond.i.i 247 %17 = load ptr addrspace(1), ptr addrspace(4) %16, align 8 248 %arrayidx.i.i.i = getelementptr inbounds i32, ptr addrspace(1) %17, i64 %ind.0.i.i 249 %18 = load i32, ptr addrspace(1) %arrayidx.i.i.i, align 4 250 %add10.i.i = add nsw i32 %18, %agg.tmp.i.sroa.9.0.copyload21 251 store i32 %add10.i.i, ptr addrspace(1) %arrayidx.i.i.i, align 4 252 %inc.i.i = add nuw i64 %ind.0.i.i, 1 253 br label %for.cond.i.i 254 255lexit1: ; preds = %for.cond.i.i, %LeaderMat.i 256 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.i 259 260wg_leader10.i: ; preds = %lexit1 261 %19 = load i32, ptr addrspace(3) @GCnt, align 4 262 %inc.i = add nsw i32 %19, 1 263 store i32 %inc.i, ptr addrspace(3) @GCnt, align 4 264 br label %wg_cf11.i 265 266wg_cf11.i: ; preds = %wg_leader10.i, %lexit1 267 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272) 268 br label %for.cond.i 269 270_ZNK11PFWGFunctorclEN4sycl3_V15groupILi1EEE.exit: ; preds = %wg_cf5.i 271 call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %agg.tmp67) 272 ret void 273} 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 nounwind 279declare 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 nounwind 294define 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 8 297 %0 = load i64, ptr %_arg_dev_ptr1, align 8 298 %1 = load i64, ptr %_arg_dev_ptr2, align 8 299 %2 = load i64, ptr %_arg_dev_ptr3, align 8 300 store i64 %2, ptr addrspace(3) @GKernel1, align 8 301 store i64 %0, ptr addrspace(3) undef, align 8 302 store i64 %1, ptr addrspace(3) undef, align 8 303 %add.ptr.i = getelementptr inbounds i32, ptr addrspace(1) %_arg_dev_ptr, i64 %2 304 store ptr addrspace(1) %add.ptr.i, ptr addrspace(3) undef, align 8 305 %3 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalSize, align 32 306 %4 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupSize, align 32 307 %5 = load i64, ptr addrspace(1) @__spirv_BuiltInNumWorkgroups, align 32 308 %6 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupId, align 32 309 call void @llvm.lifetime.start.p0(i64 32, ptr nonnull %agg.tmp67) 310 store i64 %3, ptr %agg.tmp67, align 1 311 %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 8 312 store i64 %4, ptr %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx, align 1 313 %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 16 314 store i64 %5, ptr %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx, align 1 315 %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 24 316 store i64 %6, ptr %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx, align 1 317 %7 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationIndex, align 8 318 tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272) 319 %cmpz27.i = icmp eq i64 %7, 0 320 br i1 %cmpz27.i, label %leader.i, label %merge.i 321 322leader.i: ; preds = %entry 323 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.i 325 326merge.i: ; preds = %leader.i, %entry 327 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.i 331 332wg_leader.i: ; preds = %merge.i 333 %g.ascast.i = addrspacecast ptr %agg.tmp67 to ptr addrspace(4) 334 store ptr addrspace(4) %g.ascast.i, ptr addrspace(3) @GAsCast2, align 8 335 store i32 0, ptr addrspace(3) @GCnt2, align 4 336 br label %wg_cf.i 337 338wg_cf.i: ; preds = %wg_leader.i, %merge.i 339 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272) 340 %8 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalInvocationId, align 32 341 %9 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationId, align 32 342 %cmp.i.i.i.i.i.i = icmp ult i64 %8, 2147483648 343 br label %for.cond.i 344 345for.cond.i: ; preds = %wg_cf18.i, %wg_cf.i 346 %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.i 350 351wg_leader8.i: ; preds = %for.cond.i 352 %10 = load i32, ptr addrspace(3) @GCnt2, align 4 353 %cmp.i = icmp slt i32 %10, 2 354 store i1 %cmp.i, ptr addrspace(3) @GCmp2, align 1 355 br label %wg_cf9.i 356 357wg_cf9.i: ; preds = %wg_leader8.i, %for.cond.i 358 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 1 360 br i1 %wg_val_cmp.i, label %for.body.i, label %lexit2 361 362for.body.i: ; preds = %wg_cf9.i 363 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.i 366 367TestMat25.i: ; preds = %for.body.i 368 store i64 %agg.tmp5.i.sroa.0.0, ptr addrspace(3) @WGCopy.6.0, align 8 369 store i64 ptrtoint (ptr addrspace(4) addrspacecast (ptr addrspace(3) @GKernel1 to ptr addrspace(4)) to i64), ptr addrspace(3) @WGCopy.4.0, align 8 370 store i64 5, ptr addrspace(3) @WGCopy.3.0, align 8 371 store i64 %agg.tmp4.i.sroa.0.0, ptr addrspace(3) @WGCopy.5.0, align 8 372 br label %LeaderMat22.i 373 374LeaderMat22.i: ; preds = %TestMat25.i, %for.body.i 375 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272) 376 %11 = load i64, ptr addrspace(3) @WGCopy.3.0, align 8 377 %12 = load i64, ptr addrspace(3) @WGCopy.4.0, align 8 378 %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 24 382 br label %for.cond.i.i 383 384for.cond.i.i: ; preds = %for.body.i.i, %LeaderMat22.i 385 %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, %11 387 br i1 %cmp.i.i, label %for.body.i.i, label %lexit3 388 389for.body.i.i: ; preds = %for.cond.i.i 390 call void @llvm.assume(i1 %cmp.i.i.i.i.i.i) 391 %15 = load ptr addrspace(1), ptr addrspace(4) %14, align 8 392 %arrayidx.i.i.i.i.i = getelementptr inbounds i32, ptr addrspace(1) %15, i64 %8 393 %16 = load i32, ptr addrspace(1) %arrayidx.i.i.i.i.i, align 4 394 %inc.i.i.i.i = add nsw i32 %16, 1 395 store i32 %inc.i.i.i.i, ptr addrspace(1) %arrayidx.i.i.i.i.i, align 4 396 %add.i.i = add i64 %storemerge.i.i, %4 397 br label %for.cond.i.i 398 399lexit3: ; preds = %for.cond.i.i 400 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.i 404 405TestMat.i: ; preds = %lexit3 406 store i64 ptrtoint (ptr addrspace(4) addrspacecast (ptr addrspace(3) @GKernel1 to ptr addrspace(4)) to i64), ptr addrspace(3) @WGCopy.6.0, align 8 407 store i64 %12, ptr addrspace(3) @WGCopy.4.0, align 8 408 store i64 %11, ptr addrspace(3) @WGCopy.3.0, align 8 409 store i64 2, ptr addrspace(3) @WGCopy.5.0, align 8 410 br label %LeaderMat.i 411 412LeaderMat.i: ; preds = %TestMat.i, %lexit3 413 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272) 414 %17 = load i64, ptr addrspace(3) @WGCopy.5.0, align 8 415 %18 = load i64, ptr addrspace(3) @WGCopy.6.0, align 8 416 %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 24 420 br label %for.cond.i.i19 421 422for.cond.i.i19: ; preds = %for.body.i.i22, %LeaderMat.i 423 %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, %17 425 br i1 %cmp.i.i21, label %for.body.i.i22, label %lexit4 426 427for.body.i.i22: ; preds = %for.cond.i.i19 428 call void @llvm.assume(i1 %cmp.i.i.i.i.i.i) 429 %21 = load ptr addrspace(1), ptr addrspace(4) %20, align 8 430 %arrayidx.i.i.i.i.i23 = getelementptr inbounds i32, ptr addrspace(1) %21, i64 %8 431 %22 = load i32, ptr addrspace(1) %arrayidx.i.i.i.i.i23, align 4 432 %inc.i.i.i.i25 = add nsw i32 %22, 1 433 store i32 %inc.i.i.i.i25, ptr addrspace(1) %arrayidx.i.i.i.i.i23, align 4 434 %add.i.i26 = add i64 %storemerge.i.i20, %4 435 br label %for.cond.i.i19 436 437lexit4: ; preds = %for.cond.i.i19 438 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.i 441 442wg_leader17.i: ; preds = %lexit4 443 %23 = load i32, ptr addrspace(3) @GCnt2, align 4 444 %inc.i = add nsw i32 %23, 1 445 store i32 %inc.i, ptr addrspace(3) @GCnt2, align 4 446 br label %wg_cf18.i 447 448wg_cf18.i: ; preds = %wg_leader17.i, %lexit4 449 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272) 450 br label %for.cond.i 451 452lexit2: ; preds = %wg_cf9.i 453 call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %agg.tmp67) 454 ret void 455} 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 nounwind 461define 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 8 464 %0 = load i64, ptr %_arg_dev_ptr1, align 8 465 %1 = load i64, ptr %_arg_dev_ptr2, align 8 466 %2 = load i64, ptr %_arg_dev_ptr3, align 8 467 store i64 %2, ptr addrspace(3) @GKernel2, align 8 468 store i64 %0, ptr addrspace(3) undef, align 8 469 store i64 %1, ptr addrspace(3) undef, align 8 470 %add.ptr.i = getelementptr inbounds i32, ptr addrspace(1) %_arg_dev_ptr, i64 %2 471 store ptr addrspace(1) %add.ptr.i, ptr addrspace(3) undef, align 8 472 %3 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalSize, align 32 473 %4 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupSize, align 32 474 %5 = load i64, ptr addrspace(1) @__spirv_BuiltInNumWorkgroups, align 32 475 %6 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupId, align 32 476 call void @llvm.lifetime.start.p0(i64 32, ptr nonnull %agg.tmp67) 477 store i64 %3, ptr %agg.tmp67, align 1 478 %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 8 479 store i64 %4, ptr %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx, align 1 480 %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 16 481 store i64 %5, ptr %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx, align 1 482 %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 24 483 store i64 %6, ptr %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx, align 1 484 %7 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationIndex, align 8 485 tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272) 486 %cmpz15.i = icmp eq i64 %7, 0 487 br i1 %cmpz15.i, label %leader.i, label %merge.i 488 489leader.i: ; preds = %entry 490 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.i 492 493merge.i: ; preds = %leader.i, %entry 494 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.i 498 499wg_leader.i: ; preds = %merge.i 500 %g.ascast.i = addrspacecast ptr %agg.tmp67 to ptr addrspace(4) 501 store ptr addrspace(4) %g.ascast.i, ptr addrspace(3) @GAscast3, align 8 502 store i32 0, ptr addrspace(3) @GCnt3, align 4 503 br label %wg_cf.i 504 505wg_cf.i: ; preds = %wg_leader.i, %merge.i 506 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272) 507 %8 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalInvocationId, align 32 508 %9 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationId, align 32 509 %cmp.i.i.i.i.i.i = icmp ult i64 %8, 2147483648 510 br label %for.cond.i 511 512for.cond.i: ; preds = %wg_cf11.i, %wg_cf.i 513 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.i 515 516wg_leader4.i: ; preds = %for.cond.i 517 %10 = load i32, ptr addrspace(3) @GCnt3, align 4 518 %cmp.i = icmp slt i32 %10, 2 519 store i1 %cmp.i, ptr addrspace(3) @GCmp3, align 1 520 br label %wg_cf5.i 521 522wg_cf5.i: ; preds = %wg_leader4.i, %for.cond.i 523 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 1 525 br i1 %wg_val_cmp.i, label %for.body.i, label %lexit6 526 527for.body.i: ; preds = %wg_cf5.i 528 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.i 531 532TestMat.i: ; preds = %for.body.i 533 store i64 ptrtoint (ptr addrspace(4) addrspacecast (ptr addrspace(3) @GKernel2 to ptr addrspace(4)) to i64), ptr addrspace(3) @WGCopy.10.0, align 8 534 store i64 5, ptr addrspace(3) @WGCopy.9.0, align 8 535 br label %LeaderMat.i 536 537LeaderMat.i: ; preds = %TestMat.i, %for.body.i 538 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272) 539 %11 = load i64, ptr addrspace(3) @WGCopy.9.0, align 8 540 %12 = load i64, ptr addrspace(3) @WGCopy.10.0, align 8 541 %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 24 545 br label %for.cond.i.i 546 547for.cond.i.i: ; preds = %for.body.i.i, %LeaderMat.i 548 %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, %11 550 br i1 %cmp.i.i, label %for.body.i.i, label %lexit7 551 552for.body.i.i: ; preds = %for.cond.i.i 553 %cmp5.not.i.i.i.i.i.i = icmp ne i64 %storemerge.i.i, %9 554 %cond.i.i.i.i = zext i1 %cmp5.not.i.i.i.i.i.i to i32 555 call void @llvm.assume(i1 %cmp.i.i.i.i.i.i) 556 %15 = load ptr addrspace(1), ptr addrspace(4) %14, align 8 557 %arrayidx.i.i.i.i.i = getelementptr inbounds i32, ptr addrspace(1) %15, i64 %8 558 %16 = load i32, ptr addrspace(1) %arrayidx.i.i.i.i.i, align 4 559 %add.i.i.i.i = add nsw i32 %16, %cond.i.i.i.i 560 store i32 %add.i.i.i.i, ptr addrspace(1) %arrayidx.i.i.i.i.i, align 4 561 %add.i.i = add i64 %storemerge.i.i, %4 562 br label %for.cond.i.i 563 564lexit7: ; preds = %for.cond.i.i 565 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.i 568 569wg_leader10.i: ; preds = %lexit7 570 %17 = load i32, ptr addrspace(3) @GCnt3, align 4 571 %inc.i = add nsw i32 %17, 1 572 store i32 %inc.i, ptr addrspace(3) @GCnt3, align 4 573 br label %wg_cf11.i 574 575wg_cf11.i: ; preds = %wg_leader10.i, %lexit7 576 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272) 577 br label %for.cond.i 578 579lexit6: ; preds = %wg_cf5.i 580 call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %agg.tmp67) 581 ret void 582} 583 584; Function Attrs: convergent mustprogress norecurse nounwind 585define 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 4 588 %agg.tmp67 = alloca %"class.sycl::_V1::group.15", align 8 589 %0 = load i64, ptr %_arg_dev_ptr1, align 8 590 %1 = load i64, ptr %_arg_dev_ptr2, align 8 591 %2 = load i64, ptr %_arg_dev_ptr3, align 8 592 store i64 %2, ptr addrspace(3) @GKernel3, align 8 593 store i64 %0, ptr addrspace(3) undef, align 8 594 store i64 %1, ptr addrspace(3) undef, align 8 595 %add.ptr.i = getelementptr inbounds i32, ptr addrspace(1) %_arg_dev_ptr, i64 %2 596 store ptr addrspace(1) %add.ptr.i, ptr addrspace(3) undef, align 8 597 %3 = load i64, ptr addrspace(1) undef, align 8 598 %4 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalSize, align 32 599 %5 = load i64, ptr addrspace(1) undef, align 8 600 %6 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupSize, align 32 601 %7 = load i64, ptr addrspace(1) undef, align 8 602 %8 = load i64, ptr addrspace(1) @__spirv_BuiltInNumWorkgroups, align 32 603 %9 = load i64, ptr addrspace(1) undef, align 8 604 %10 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupId, align 32 605 call void @llvm.lifetime.start.p0(i64 64, ptr nonnull %agg.tmp67) 606 store i64 %3, ptr %agg.tmp67, align 1 607 %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 8 608 store i64 %4, ptr %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx, align 1 609 %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 16 610 store i64 %5, ptr %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx, align 1 611 %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 24 612 store i64 %6, ptr %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx, align 1 613 %agg.tmp6.sroa.5.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 32 614 store i64 %7, ptr %agg.tmp6.sroa.5.0.agg.tmp67.sroa_idx, align 1 615 %agg.tmp6.sroa.6.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 40 616 store i64 %8, ptr %agg.tmp6.sroa.6.0.agg.tmp67.sroa_idx, align 1 617 %agg.tmp6.sroa.7.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 48 618 store i64 %9, ptr %agg.tmp6.sroa.7.0.agg.tmp67.sroa_idx, align 1 619 %agg.tmp6.sroa.8.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 56 620 store i64 %10, ptr %agg.tmp6.sroa.8.0.agg.tmp67.sroa_idx, align 1 621 %11 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationIndex, align 8 622 tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272) 623 %cmpz32.i = icmp eq i64 %11, 0 624 br i1 %cmpz32.i, label %leader.i, label %merge.i 625 626leader.i: ; preds = %entry 627 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.i 629 630merge.i: ; preds = %leader.i, %entry 631 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.i 636 637wg_leader.i: ; preds = %merge.i 638 %g.ascast.i = addrspacecast ptr %agg.tmp67 to ptr addrspace(4) 639 store ptr addrspace(4) %g.ascast.i, ptr addrspace(3) @GAsCast4, align 8 640 call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %priv.i) 641 store i32 0, ptr addrspace(3) @GCnt4, align 4 642 br label %wg_cf.i 643 644wg_cf.i: ; preds = %wg_leader.i, %merge.i 645 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272) 646 %12 = load i64, ptr addrspace(1) undef, align 8 647 %13 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalInvocationId, align 32 648 %14 = load i64, ptr addrspace(1) undef, align 8 649 %15 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationId, align 32 650 %mul.i.i.i.i.i.i = mul i64 %12, %4 651 %add.i.i.i.i.i.i = add i64 %mul.i.i.i.i.i.i, %13 652 %cmp.i.i.i.i.i.i = icmp ult i64 %add.i.i.i.i.i.i, 2147483648 653 %conv.i.i.i.i.i = trunc i64 %add.i.i.i.i.i.i to i32 654 %y.i.i.i.i.i = getelementptr inbounds i8, ptr %priv.i, i64 4 655 br label %for.cond.i 656 657for.cond.i: ; preds = %wg_cf20.i, %wg_cf.i 658 %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.i 667 668wg_leader10.i: ; preds = %for.cond.i 669 %16 = load i32, ptr addrspace(3) @GCnt4, align 4 670 %cmp.i = icmp slt i32 %16, 2 671 store i1 %cmp.i, ptr addrspace(3) @GCmp4, align 1 672 br label %wg_cf11.i 673 674wg_cf11.i: ; preds = %wg_leader10.i, %for.cond.i 675 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 1 677 br i1 %wg_val_cmp.i, label %for.body.i, label %for.end.i 678 679for.body.i: ; preds = %wg_cf11.i 680 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.i 682 683wg_leader13.i: ; preds = %for.body.i 684 br label %wg_cf14.i 685 686wg_cf14.i: ; preds = %wg_leader13.i, %for.body.i 687 %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.i 693 694TestMat30.i: ; preds = %wg_cf14.i 695 store i64 %agg.tmp.i.sroa.0.1, ptr addrspace(3) @WGCopy.13.0, align 8 696 store i64 %agg.tmp.i.sroa.8.1, ptr addrspace(3) @WGCopy.13.1, align 8 697 store ptr addrspace(4) %agg.tmp2.i.sroa.0.1, ptr addrspace(3) @WGCopy.14.0, align 8 698 store ptr addrspace(4) %agg.tmp2.i.sroa.8.1, ptr addrspace(3) @WGCopy.14.1, align 8 699 store i64 %agg.tmp5.i.sroa.0.0, ptr addrspace(3) @WGCopy.15.0, align 8 700 store i64 %agg.tmp5.i.sroa.8.0, ptr addrspace(3) @WGCopy.15.1, align 8 701 store ptr addrspace(4) %priv.ascast.i, ptr addrspace(3) @WGCopy.16.0, align 8 702 store ptr addrspace(4) %agg.tmp6.i.sroa.9.0, ptr addrspace(3) @WGCopy.16.1, align 8 703 br label %LeaderMat27.i 704 705LeaderMat27.i: ; preds = %TestMat30.i, %wg_cf14.i 706 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 8 708 %agg.tmp6.i.sroa.9.0.copyload = load ptr addrspace(4), ptr addrspace(3) @WGCopy.16.1, align 8 709 %agg.tmp5.i.sroa.0.0.copyload = load i64, ptr addrspace(3) @WGCopy.15.0, align 8 710 %agg.tmp5.i.sroa.8.0.copyload = load i64, ptr addrspace(3) @WGCopy.15.1, align 8 711 %agg.tmp2.i.sroa.0.0.copyload = load ptr addrspace(4), ptr addrspace(3) @WGCopy.14.0, align 8 712 %agg.tmp.i.sroa.0.0.copyload = load i64, ptr addrspace(3) @WGCopy.13.0, align 8 713 %agg.tmp.i.sroa.8.0.copyload = load i64, ptr addrspace(3) @WGCopy.13.1, align 8 714 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 24 717 br label %for.cond.i.i 718 719for.cond.i.i: ; preds = %lexit10, %LeaderMat27.i 720 %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.copyload 722 br i1 %cmp.i.i, label %for.cond.i.i.i, label %lexit11 723 724for.cond.i.i.i: ; preds = %for.body.i.i.i, %for.cond.i.i 725 %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.copyload 727 br i1 %cmp.i.i.i, label %for.body.i.i.i, label %lexit10 728 729for.body.i.i.i: ; preds = %for.cond.i.i.i 730 call void @llvm.assume(i1 %cmp.i.i.i.i.i.i) 731 %18 = load ptr addrspace(1), ptr addrspace(4) %17, align 8 732 %arrayidx.i.i.i.i.i.i = getelementptr inbounds i32, ptr addrspace(1) %18, i64 %add.i.i.i.i.i.i 733 %19 = load i32, ptr addrspace(1) %arrayidx.i.i.i.i.i.i, align 4 734 %inc.i.i.i.i.i = add nsw i32 %19, 1 735 store i32 %inc.i.i.i.i.i, ptr addrspace(1) %arrayidx.i.i.i.i.i.i, align 4 736 store i32 %conv.i.i.i.i.i, ptr %priv.i, align 4 737 store i32 5, ptr %y.i.i.i.i.i, align 4 738 %add.i.i.i = add i64 %storemerge.i.i.i, %6 739 br label %for.cond.i.i.i 740 741lexit10: ; preds = %for.cond.i.i.i 742 %add.i.i = add i64 %storemerge.i.i, %5 743 br label %for.cond.i.i 744 745lexit11: ; preds = %for.cond.i.i 746 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.i 749 750wg_leader16.i: ; preds = %lexit11 751 br label %wg_cf17.i 752 753wg_cf17.i: ; preds = %wg_leader16.i, %lexit11 754 %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.i 760 761TestMat.i: ; preds = %wg_cf17.i 762 store i64 %agg.tmp.i.sroa.0.0.copyload, ptr addrspace(3) @WGCopy.13.0, align 8 763 store i64 %agg.tmp.i.sroa.8.0.copyload, ptr addrspace(3) @WGCopy.13.1, align 8 764 store ptr addrspace(4) %agg.tmp2.i.sroa.0.0.copyload, ptr addrspace(3) @WGCopy.14.0, align 8 765 store ptr addrspace(4) %priv.ascast.i, ptr addrspace(3) @WGCopy.14.1, align 8 766 store i64 %agg.tmp5.i.sroa.0.1, ptr addrspace(3) @WGCopy.15.0, align 8 767 store i64 %agg.tmp5.i.sroa.8.1, ptr addrspace(3) @WGCopy.15.1, align 8 768 store ptr addrspace(4) %agg.tmp6.i.sroa.0.1, ptr addrspace(3) @WGCopy.16.0, align 8 769 store ptr addrspace(4) %agg.tmp6.i.sroa.9.1, ptr addrspace(3) @WGCopy.16.1, align 8 770 br label %LeaderMat.i 771 772LeaderMat.i: ; preds = %TestMat.i, %wg_cf17.i 773 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 8 775 %agg.tmp5.i.sroa.0.0.copyload44 = load i64, ptr addrspace(3) @WGCopy.15.0, align 8 776 %agg.tmp5.i.sroa.8.0.copyload48 = load i64, ptr addrspace(3) @WGCopy.15.1, align 8 777 %agg.tmp2.i.sroa.0.0.copyload52 = load ptr addrspace(4), ptr addrspace(3) @WGCopy.14.0, align 8 778 %agg.tmp2.i.sroa.8.0.copyload56 = load ptr addrspace(4), ptr addrspace(3) @WGCopy.14.1, align 8 779 %agg.tmp.i.sroa.0.0.copyload60 = load i64, ptr addrspace(3) @WGCopy.13.0, align 8 780 %agg.tmp.i.sroa.8.0.copyload64 = load i64, ptr addrspace(3) @WGCopy.13.1, align 8 781 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 24 784 br label %for.cond.i.i25 785 786for.cond.i.i25: ; preds = %lexit12, %LeaderMat.i 787 %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.copyload44 789 br i1 %cmp.i.i27, label %for.cond.i.i.i28, label %lexit13 790 791for.cond.i.i.i28: ; preds = %for.body.i.i.i32, %for.cond.i.i25 792 %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.copyload48 794 br i1 %cmp.i.i.i30, label %for.body.i.i.i32, label %lexit12 795 796for.body.i.i.i32: ; preds = %for.cond.i.i.i28 797 %21 = load i32, ptr %priv.i, align 4 798 %22 = load i32, ptr %y.i.i.i.i.i, align 4 799 %add.i.i.i.i.i = add nsw i32 %21, %22 800 call void @llvm.assume(i1 %cmp.i.i.i.i.i.i) 801 %23 = load ptr addrspace(1), ptr addrspace(4) %20, align 8 802 %arrayidx.i.i.i.i.i.i33 = getelementptr inbounds i32, ptr addrspace(1) %23, i64 %add.i.i.i.i.i.i 803 %24 = load i32, ptr addrspace(1) %arrayidx.i.i.i.i.i.i33, align 4 804 %add4.i.i.i.i.i = add nsw i32 %24, %add.i.i.i.i.i 805 store i32 %add4.i.i.i.i.i, ptr addrspace(1) %arrayidx.i.i.i.i.i.i33, align 4 806 %add.i.i.i35 = add i64 %storemerge.i.i.i29, %6 807 br label %for.cond.i.i.i28 808 809lexit12: ; preds = %for.cond.i.i.i28 810 %add.i.i31 = add i64 %storemerge.i.i26, %5 811 br label %for.cond.i.i25 812 813lexit13: ; preds = %for.cond.i.i25 814 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.i 817 818wg_leader19.i: ; preds = %lexit13 819 %25 = load i32, ptr addrspace(3) @GCnt4, align 4 820 %inc.i = add nsw i32 %25, 1 821 store i32 %inc.i, ptr addrspace(3) @GCnt4, align 4 822 br label %wg_cf20.i 823 824wg_cf20.i: ; preds = %wg_leader19.i, %lexit13 825 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272) 826 br label %for.cond.i 827 828for.end.i: ; preds = %wg_cf11.i 829 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272) 830 br i1 %cmpz32.i, label %wg_leader22.i, label %lexit14 831 832wg_leader22.i: ; preds = %for.end.i 833 call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %priv.i) 834 br label %lexit14 835 836lexit14: ; preds = %wg_leader22.i, %for.end.i 837 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 void 840} 841 842; Function Attrs: convergent mustprogress norecurse nounwind 843define 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 8 846 %0 = load i64, ptr %_arg_dev_ptr1, align 8 847 %1 = load i64, ptr %_arg_dev_ptr2, align 8 848 %2 = load i64, ptr %_arg_dev_ptr3, align 8 849 store i64 %2, ptr addrspace(3) @GKernel4, align 8 850 store i64 %0, ptr addrspace(3) undef, align 8 851 store i64 %1, ptr addrspace(3) undef, align 8 852 %add.ptr.i = getelementptr inbounds i32, ptr addrspace(1) %_arg_dev_ptr, i64 %2 853 store ptr addrspace(1) %add.ptr.i, ptr addrspace(3) undef, align 8 854 %3 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalSize, align 32 855 %4 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupSize, align 32 856 %5 = load i64, ptr addrspace(1) @__spirv_BuiltInNumWorkgroups, align 32 857 %6 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupId, align 32 858 call void @llvm.lifetime.start.p0(i64 32, ptr nonnull %agg.tmp67) 859 store i64 %3, ptr %agg.tmp67, align 1 860 %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 8 861 store i64 %4, ptr %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx, align 1 862 %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 16 863 store i64 %5, ptr %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx, align 1 864 %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 24 865 store i64 %6, ptr %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx, align 1 866 %7 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationIndex, align 8 867 tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272) 868 %cmpz16.i = icmp eq i64 %7, 0 869 br i1 %cmpz16.i, label %leader.i, label %merge.i 870 871leader.i: ; preds = %entry 872 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.i 874 875merge.i: ; preds = %leader.i, %entry 876 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.i 880 881wg_leader.i: ; preds = %merge.i 882 %g.ascast.i = addrspacecast ptr %agg.tmp67 to ptr addrspace(4) 883 store ptr addrspace(4) %g.ascast.i, ptr addrspace(3) @GAsCast5, align 8 884 store i32 0, ptr addrspace(3) @GCnt5, align 4 885 br label %wg_cf.i 886 887wg_cf.i: ; preds = %wg_leader.i, %merge.i 888 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 8 890 %8 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationId, align 32 891 %9 = trunc i64 %4 to i32 892 br label %for.cond.i 893 894for.cond.i: ; preds = %wg_cf12.i, %wg_cf.i 895 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.i 897 898wg_leader5.i: ; preds = %for.cond.i 899 %10 = load i32, ptr addrspace(3) @GCnt5, align 4 900 %cmp.i = icmp slt i32 %10, 2 901 store i1 %cmp.i, ptr addrspace(3) @GCmp5, align 1 902 br label %wg_cf6.i 903 904wg_cf6.i: ; preds = %wg_leader5.i, %for.cond.i 905 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 1 907 br i1 %wg_val_cmp.i, label %for.body.i, label %lexit20 908 909for.body.i: ; preds = %wg_cf6.i 910 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.i 913 914TestMat.i: ; preds = %for.body.i 915 store ptr addrspace(4) %wg_val_g.ascast.i, ptr addrspace(3) @WGCopy.20.0, align 8 916 store ptr addrspace(4) addrspacecast (ptr addrspace(3) @GKernel4 to ptr addrspace(4)), ptr addrspace(3) @WGCopy.20.1, align 8 917 store i64 5, ptr addrspace(3) @WGCopy.19.0, align 8 918 br label %LeaderMat.i 919 920LeaderMat.i: ; preds = %TestMat.i, %for.body.i 921 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272) 922 %11 = load i64, ptr addrspace(3) @WGCopy.19.0, align 8 923 %agg.tmp2.i.sroa.0.0.copyload = load ptr addrspace(4), ptr addrspace(3) @WGCopy.20.0, align 8 924 %agg.tmp2.i.sroa.6.0.copyload = load ptr addrspace(4), ptr addrspace(3) @WGCopy.20.1, align 8 925 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 24 928 %12 = getelementptr inbounds i8, ptr addrspace(4) %agg.tmp2.i.sroa.6.0.copyload, i64 24 929 %13 = trunc i64 %11 to i32 930 br label %for.cond.i.i 931 932for.cond.i.i: ; preds = %for.body.i.i, %LeaderMat.i 933 %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, %11 935 br i1 %cmp.i.i, label %for.body.i.i, label %lexit21 936 937for.body.i.i: ; preds = %for.cond.i.i 938 %14 = load i64, ptr addrspace(4) %index.i.i.i.i.i, align 8 939 %mul.i.i.i.i = mul i64 %14, 10 940 %mul3.i.i.i.i = shl i64 %storemerge.i.i, 1 941 %add.i.i.i.i = add i64 %mul.i.i.i.i, %mul3.i.i.i.i 942 %15 = load ptr addrspace(1), ptr addrspace(4) %12, align 8 943 %arrayidx.i.i.i.i.i = getelementptr inbounds i32, ptr addrspace(1) %15, i64 %add.i.i.i.i 944 %16 = load i32, ptr addrspace(1) %arrayidx.i.i.i.i.i, align 4 945 %conv9.i.i.i.i = add i32 %16, %13 946 store i32 %conv9.i.i.i.i, ptr addrspace(1) %arrayidx.i.i.i.i.i, align 4 947 %add14.i.i.i.i = or disjoint i64 %add.i.i.i.i, 1 948 %17 = load ptr addrspace(1), ptr addrspace(4) %12, align 8 949 %arrayidx.i25.i.i.i.i = getelementptr inbounds i32, ptr addrspace(1) %17, i64 %add14.i.i.i.i 950 %18 = load i32, ptr addrspace(1) %arrayidx.i25.i.i.i.i, align 4 951 %conv18.i.i.i.i = add i32 %18, %9 952 store i32 %conv18.i.i.i.i, ptr addrspace(1) %arrayidx.i25.i.i.i.i, align 4 953 %add.i.i = add i64 %storemerge.i.i, %4 954 br label %for.cond.i.i 955 956lexit21: ; preds = %for.cond.i.i 957 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.i 960 961wg_leader11.i: ; preds = %lexit21 962 %19 = load i32, ptr addrspace(3) @GCnt5, align 4 963 %inc.i = add nsw i32 %19, 1 964 store i32 %inc.i, ptr addrspace(3) @GCnt5, align 4 965 br label %wg_cf12.i 966 967wg_cf12.i: ; preds = %wg_leader11.i, %lexit21 968 call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272) 969 br label %for.cond.i 970 971lexit20: ; preds = %wg_cf6.i 972 call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %agg.tmp67) 973 ret void 974} 975