; This is an excerpt from the SYCL end-to-end test suite, cleaned out from unrelevant details,
; that reproduced multiple cases of the issues when OpPhi's result type mismatches with operand types.
; The only pass criterion is that spirv-val considers output valid.

; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val --target-env spv1.4 %}

%struct.PFWGFunctor = type { i64, i64, i32, i32, %"class.sycl::_V1::accessor" }
%"class.sycl::_V1::accessor" = type { %"class.sycl::_V1::detail::AccessorImplDevice", %union.anon }
%"class.sycl::_V1::detail::AccessorImplDevice" = type { %"class.sycl::_V1::range", %"class.sycl::_V1::range", %"class.sycl::_V1::range" }
%"class.sycl::_V1::range" = type { %"class.sycl::_V1::detail::array" }
%"class.sycl::_V1::detail::array" = type { [1 x i64] }
%union.anon = type { ptr addrspace(1) }
%class.anon.2 = type { %"class.sycl::_V1::accessor" }
%"class.sycl::_V1::group" = type { %"class.sycl::_V1::range", %"class.sycl::_V1::range", %"class.sycl::_V1::range", %"class.sycl::_V1::range" }
%"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" }
%"class.sycl::_V1::range.16" = type { %"class.sycl::_V1::detail::array.17" }
%"class.sycl::_V1::detail::array.17" = type { [2 x i64] }
%"class.sycl::_V1::private_memory" = type { %struct.MyStruct }
%struct.MyStruct = type { i32, i32 }

@GFunctor = internal addrspace(3) global %struct.PFWGFunctor undef, align 8
@WI.0 = internal unnamed_addr addrspace(3) global i64 undef, align 8
@WI.1 = internal unnamed_addr addrspace(3) global i64 undef, align 8
@WI.2 = internal unnamed_addr addrspace(3) global i64 undef, align 8
@WI.3 = internal unnamed_addr addrspace(3) global i64 undef, align 8
@WI.4 = internal unnamed_addr addrspace(3) global i32 undef, align 8
@WI.6 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 8
@GCnt = internal unnamed_addr addrspace(3) global i32 undef, align 4
@__spirv_BuiltInNumWorkgroups = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
@GKernel1 = internal addrspace(3) global %class.anon.2 undef, align 8
@GCnt2 = internal unnamed_addr addrspace(3) global i32 undef, align 4
@GKernel2 = internal addrspace(3) global %class.anon.2 undef, align 8
@GCnt3 = internal unnamed_addr addrspace(3) global i32 undef, align 4
@GKernel3 = internal addrspace(3) global %class.anon.2 undef, align 8
@GCnt4 = internal unnamed_addr addrspace(3) global i32 undef, align 4
@GKernel4 = internal addrspace(3) global %class.anon.2 undef, align 8
@GCnt5 = internal unnamed_addr addrspace(3) global i32 undef, align 4
@__spirv_BuiltInLocalInvocationIndex = external local_unnamed_addr addrspace(1) constant i64, align 8
@GThis = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 8
@GAsCast = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 8
@GCmp = internal unnamed_addr addrspace(3) global i1 undef, align 1
@WGCopy = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 8
@WGCopy.1.0 = internal unnamed_addr addrspace(3) global i64 undef, align 16
@WGCopy.1.1 = internal unnamed_addr addrspace(3) global i64 undef, align 16
@WGCopy.1.2 = internal unnamed_addr addrspace(3) global i64 undef, align 16
@WGCopy.1.3 = internal unnamed_addr addrspace(3) global i64 undef, align 16
@WGCopy.1.4 = internal unnamed_addr addrspace(3) global i32 undef, align 16
@WGCopy.1.5 = internal unnamed_addr addrspace(3) global i32 undef, align 16
@WGCopy.1.6 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 16
@ArgShadow = internal unnamed_addr addrspace(3) global %"class.sycl::_V1::group" undef, align 16
@GAsCast2 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 8
@GCmp2 = internal unnamed_addr addrspace(3) global i1 undef, align 1
@WGCopy.3.0 = internal unnamed_addr addrspace(3) global i64 undef, align 8
@WGCopy.4.0 = internal unnamed_addr addrspace(3) global i64 undef, align 8
@WGCopy.5.0 = internal unnamed_addr addrspace(3) global i64 undef, align 8
@WGCopy.6.0 = internal unnamed_addr addrspace(3) global i64 undef, align 8
@ArgShadow.7 = internal unnamed_addr addrspace(3) global %"class.sycl::_V1::group" undef, align 16
@GAscast3 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 8
@GCmp3 = internal unnamed_addr addrspace(3) global i1 undef, align 1
@WGCopy.9.0 = internal unnamed_addr addrspace(3) global i64 undef, align 8
@WGCopy.10.0 = internal unnamed_addr addrspace(3) global i64 undef, align 8
@ArgShadow.11 = internal unnamed_addr addrspace(3) global %"class.sycl::_V1::group" undef, align 16
@GAsCast4 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 8
@GCmp4 = internal unnamed_addr addrspace(3) global i1 undef, align 1
@WGCopy.13.0 = internal unnamed_addr addrspace(3) global i64 undef, align 8
@WGCopy.13.1 = internal unnamed_addr addrspace(3) global i64 undef, align 8
@WGCopy.14.0 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 8
@WGCopy.14.1 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 8
@WGCopy.15.0 = internal unnamed_addr addrspace(3) global i64 undef, align 8
@WGCopy.15.1 = internal unnamed_addr addrspace(3) global i64 undef, align 8
@WGCopy.16.0 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 8
@WGCopy.16.1 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 8
@ArgShadow.17 = internal unnamed_addr addrspace(3) global %"class.sycl::_V1::group.15" undef, align 16
@GAsCast5 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 8
@GCmp5 = internal unnamed_addr addrspace(3) global i1 undef, align 1
@WGCopy.19.0 = internal unnamed_addr addrspace(3) global i64 undef, align 8
@WGCopy.20.0 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 8
@WGCopy.20.1 = internal unnamed_addr addrspace(3) global ptr addrspace(4) undef, align 8
@ArgShadow.21 = internal unnamed_addr addrspace(3) global %"class.sycl::_V1::group" undef, align 16
@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
@__spirv_BuiltInGlobalSize = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
@__spirv_BuiltInLocalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
@__spirv_BuiltInWorkgroupId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
@__spirv_BuiltInWorkgroupSize = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32

; Function Attrs: convergent mustprogress norecurse nounwind
define 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) {
entry:
  %agg.tmp67 = alloca %"class.sycl::_V1::group", align 8
  store i64 %_arg_wg_chunk, ptr addrspace(3) @GFunctor, align 8
  store i64 %_arg_range_length, ptr addrspace(3) undef, align 8
  store i32 %_arg_n_iter, ptr addrspace(3) undef, align 8
  store i32 %_arg_addend, ptr addrspace(3) undef, align 4
  %0 = load i64, ptr %_arg_dev_ptr1, align 8
  %1 = load i64, ptr %_arg_dev_ptr2, align 8
  %2 = load i64, ptr %_arg_dev_ptr3, align 8
  store i64 %2, ptr addrspace(3) undef, align 8
  store i64 %0, ptr addrspace(3) undef, align 8
  store i64 %1, ptr addrspace(3) undef, align 8
  %add.ptr.i = getelementptr inbounds i32, ptr addrspace(1) %_arg_dev_ptr, i64 %2
  store ptr addrspace(1) %add.ptr.i, ptr addrspace(3) undef, align 8
  %3 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalSize, align 32
  %4 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupSize, align 32
  %5 = load i64, ptr addrspace(1) @__spirv_BuiltInNumWorkgroups, align 32
  %6 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupId, align 32
  call void @llvm.lifetime.start.p0(i64 32, ptr nonnull %agg.tmp67)
  store i64 %3, ptr %agg.tmp67, align 1
  %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 8
  store i64 %4, ptr %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx, align 1
  %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 16
  store i64 %5, ptr %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx, align 1
  %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 24
  store i64 %6, ptr %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx, align 1
  %7 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationIndex, align 8
  tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  %cmpz15.i = icmp eq i64 %7, 0
  br i1 %cmpz15.i, label %leader.i, label %merge.i

leader.i:                                         ; preds = %entry
  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)
  br label %merge.i

merge.i:                                          ; preds = %leader.i, %entry
  tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  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)
  tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br i1 %cmpz15.i, label %wg_leader.i, label %wg_cf.i

wg_leader.i:                                      ; preds = %merge.i
  %g.ascast.i = addrspacecast ptr %agg.tmp67 to ptr addrspace(4)
  store ptr addrspace(4) %g.ascast.i, ptr addrspace(3) @GAsCast, align 8
  store ptr addrspace(4) addrspacecast (ptr addrspace(3) @GFunctor to ptr addrspace(4)), ptr addrspace(3) @GThis, align 8
  %8 = load i32, ptr addrspace(3) undef, align 4
  %9 = load i64, ptr addrspace(3) @GFunctor, align 8
  %index.i = getelementptr inbounds i8, ptr %agg.tmp67, i64 24
  %10 = load i64, ptr %index.i, align 8
  %mul.i = mul i64 %9, %10
  %localRange.i = getelementptr inbounds i8, ptr %agg.tmp67, i64 8
  %11 = load i64, ptr %localRange.i, align 8
  %12 = load i64, ptr addrspace(3) undef, align 8
  store i64 %9, ptr addrspace(3) @WI.0, align 8
  store i64 %11, ptr addrspace(3) @WI.1, align 8
  store i64 %mul.i, ptr addrspace(3) @WI.2, align 8
  store i64 %12, ptr addrspace(3) @WI.3, align 8
  store i32 %8, ptr addrspace(3) @WI.4, align 8
  store ptr addrspace(4) undef, ptr addrspace(3) @WI.6, align 8
  store i32 0, ptr addrspace(3) @GCnt, align 4
  br label %wg_cf.i

wg_cf.i:                                          ; preds = %wg_leader.i, %merge.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  %wg_val_this1.i = load ptr addrspace(4), ptr addrspace(3) @GThis, align 8
  %n_iter.i = getelementptr inbounds i8, ptr addrspace(4) %wg_val_this1.i, i64 16
  %13 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationId, align 32
  br label %for.cond.i

for.cond.i:                                       ; preds = %wg_cf11.i, %wg_cf.i
  %agg.tmp.i.sroa.0.0 = phi i64 [ undef, %wg_cf.i ], [ %agg.tmp.i.sroa.0.0.copyload13, %wg_cf11.i ]
  %agg.tmp.i.sroa.6.0 = phi i64 [ undef, %wg_cf.i ], [ %agg.tmp.i.sroa.6.0.copyload15, %wg_cf11.i ]
  %agg.tmp.i.sroa.7.0 = phi i64 [ undef, %wg_cf.i ], [ %agg.tmp.i.sroa.7.0.copyload17, %wg_cf11.i ]
  %agg.tmp.i.sroa.8.0 = phi i64 [ undef, %wg_cf.i ], [ %agg.tmp.i.sroa.8.0.copyload19, %wg_cf11.i ]
  %agg.tmp.i.sroa.9.0 = phi i32 [ undef, %wg_cf.i ], [ %agg.tmp.i.sroa.9.0.copyload21, %wg_cf11.i ]
  %agg.tmp.i.sroa.10.0 = phi i32 [ undef, %wg_cf.i ], [ %agg.tmp.i.sroa.10.0.copyload23, %wg_cf11.i ]
  %agg.tmp.i.sroa.11.0 = phi ptr addrspace(4) [ undef, %wg_cf.i ], [ %agg.tmp.i.sroa.11.0.copyload25, %wg_cf11.i ]
  %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 ]
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br i1 %cmpz15.i, label %wg_leader4.i, label %wg_cf5.i

wg_leader4.i:                                     ; preds = %for.cond.i
  %14 = load i32, ptr addrspace(3) @GCnt, align 4
  %15 = load i32, ptr addrspace(4) %n_iter.i, align 8
  %cmp.i = icmp slt i32 %14, %15
  store i1 %cmp.i, ptr addrspace(3) @GCmp, align 1
  br label %wg_cf5.i

wg_cf5.i:                                         ; preds = %wg_leader4.i, %for.cond.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  %wg_val_cmp.i = load i1, ptr addrspace(3) @GCmp, align 1
  br i1 %wg_val_cmp.i, label %for.body.i, label %_ZNK11PFWGFunctorclEN4sycl3_V15groupILi1EEE.exit

for.body.i:                                       ; preds = %wg_cf5.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br i1 %cmpz15.i, label %wg_leader7.i, label %wg_cf8.i

wg_leader7.i:                                     ; preds = %for.body.i
  %agg.tmp.i.sroa.0.0.copyload = load i64, ptr addrspace(3) @WI.0, align 8
  %agg.tmp.i.sroa.6.0.copyload = load i64, ptr addrspace(3) @WI.1, align 8
  %agg.tmp.i.sroa.7.0.copyload = load i64, ptr addrspace(3) @WI.2, align 8
  %agg.tmp.i.sroa.8.0.copyload = load i64, ptr addrspace(3) @WI.3, align 8
  %agg.tmp.i.sroa.9.0.copyload = load i32, ptr addrspace(3) @WI.4, align 8
  %agg.tmp.i.sroa.11.0.copyload = load ptr addrspace(4), ptr addrspace(3) @WI.6, align 8
  br label %wg_cf8.i

wg_cf8.i:                                         ; preds = %wg_leader7.i, %for.body.i
  %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 ]
  %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 ]
  %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 ]
  %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 ]
  %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 ]
  %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 ]
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br i1 %cmpz15.i, label %TestMat.i, label %LeaderMat.i

TestMat.i:                                        ; preds = %wg_cf8.i
  store i64 %agg.tmp.i.sroa.0.1, ptr addrspace(3) @WGCopy.1.0, align 16
  store i64 %agg.tmp.i.sroa.6.1, ptr addrspace(3) @WGCopy.1.1, align 16
  store i64 %agg.tmp.i.sroa.7.1, ptr addrspace(3) @WGCopy.1.2, align 16
  store i64 %agg.tmp.i.sroa.8.1, ptr addrspace(3) @WGCopy.1.3, align 16
  store i32 %agg.tmp.i.sroa.9.1, ptr addrspace(3) @WGCopy.1.4, align 16
  store i32 %agg.tmp.i.sroa.10.0, ptr addrspace(3) @WGCopy.1.5, align 16
  store ptr addrspace(4) %agg.tmp.i.sroa.11.1, ptr addrspace(3) @WGCopy.1.6, align 16
  store ptr addrspace(4) %this.addr.0.i, ptr addrspace(3) @WGCopy, align 8
  br label %LeaderMat.i

LeaderMat.i:                                      ; preds = %TestMat.i, %wg_cf8.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  %mat_ld13.i = load ptr addrspace(4), ptr addrspace(3) @WGCopy, align 8
  %agg.tmp.i.sroa.0.0.copyload13 = load i64, ptr addrspace(3) @WGCopy.1.0, align 16
  %agg.tmp.i.sroa.6.0.copyload15 = load i64, ptr addrspace(3) @WGCopy.1.1, align 16
  %agg.tmp.i.sroa.7.0.copyload17 = load i64, ptr addrspace(3) @WGCopy.1.2, align 16
  %agg.tmp.i.sroa.8.0.copyload19 = load i64, ptr addrspace(3) @WGCopy.1.3, align 16
  %agg.tmp.i.sroa.9.0.copyload21 = load i32, ptr addrspace(3) @WGCopy.1.4, align 16
  %agg.tmp.i.sroa.10.0.copyload23 = load i32, ptr addrspace(3) @WGCopy.1.5, align 16
  %agg.tmp.i.sroa.11.0.copyload25 = load ptr addrspace(4), ptr addrspace(3) @WGCopy.1.6, align 16
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef 2, i32 noundef 2, i32 noundef 272)
  %cmp.not.i.i = icmp ult i64 %13, %agg.tmp.i.sroa.0.0.copyload13
  br i1 %cmp.not.i.i, label %if.end.i.i, label %lexit1

if.end.i.i:                                       ; preds = %LeaderMat.i
  %add.i.i = add i64 %agg.tmp.i.sroa.0.0.copyload13, %agg.tmp.i.sroa.6.0.copyload15
  %sub.i.i = add i64 %add.i.i, -1
  %div.i.i = udiv i64 %sub.i.i, %agg.tmp.i.sroa.6.0.copyload15
  %mul.i.i = mul i64 %13, %div.i.i
  %add4.i.i = add i64 %agg.tmp.i.sroa.7.0.copyload17, %mul.i.i
  %add6.i.i = add i64 %add4.i.i, %div.i.i
  %.sroa.speculated.i.i = call i64 @llvm.umin.i64(i64 %agg.tmp.i.sroa.8.0.copyload19, i64 %add6.i.i)
  %16 = getelementptr inbounds i8, ptr addrspace(4) %agg.tmp.i.sroa.11.0.copyload25, i64 24
  br label %for.cond.i.i

for.cond.i.i:                                     ; preds = %for.body.i.i, %if.end.i.i
  %ind.0.i.i = phi i64 [ %add4.i.i, %if.end.i.i ], [ %inc.i.i, %for.body.i.i ]
  %cmp8.i.i = icmp ult i64 %ind.0.i.i, %.sroa.speculated.i.i
  br i1 %cmp8.i.i, label %for.body.i.i, label %lexit1

for.body.i.i:                                     ; preds = %for.cond.i.i
  %17 = load ptr addrspace(1), ptr addrspace(4) %16, align 8
  %arrayidx.i.i.i = getelementptr inbounds i32, ptr addrspace(1) %17, i64 %ind.0.i.i
  %18 = load i32, ptr addrspace(1) %arrayidx.i.i.i, align 4
  %add10.i.i = add nsw i32 %18, %agg.tmp.i.sroa.9.0.copyload21
  store i32 %add10.i.i, ptr addrspace(1) %arrayidx.i.i.i, align 4
  %inc.i.i = add nuw i64 %ind.0.i.i, 1
  br label %for.cond.i.i

lexit1: ; preds = %for.cond.i.i, %LeaderMat.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef 2, i32 noundef 2, i32 noundef 272)
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br i1 %cmpz15.i, label %wg_leader10.i, label %wg_cf11.i

wg_leader10.i:                                    ; preds = %lexit1
  %19 = load i32, ptr addrspace(3) @GCnt, align 4
  %inc.i = add nsw i32 %19, 1
  store i32 %inc.i, ptr addrspace(3) @GCnt, align 4
  br label %wg_cf11.i

wg_cf11.i:                                        ; preds = %wg_leader10.i, %lexit1
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br label %for.cond.i

_ZNK11PFWGFunctorclEN4sycl3_V15groupILi1EEE.exit: ; preds = %wg_cf5.i
  call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %agg.tmp67)
  ret void
}

; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite)
declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture)

; Function Attrs: convergent nounwind
declare dso_local spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef, i32 noundef, i32 noundef)

; Function Attrs: nocallback nofree nounwind willreturn memory(argmem: readwrite)
declare void @llvm.memcpy.p3.p0.i64(ptr addrspace(3) noalias nocapture writeonly, ptr noalias nocapture readonly, i64, i1 immarg)

; Function Attrs: nocallback nofree nounwind willreturn memory(argmem: readwrite)
declare void @llvm.memcpy.p0.p3.i64(ptr noalias nocapture writeonly, ptr addrspace(3) noalias nocapture readonly, i64, i1 immarg)

; Function Attrs: nocallback nofree nosync nounwind speculatable willreturn memory(none)
declare i64 @llvm.umin.i64(i64, i64)

; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite)
declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture)

; Function Attrs: convergent mustprogress norecurse nounwind
define 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) {
entry:
  %agg.tmp67 = alloca %"class.sycl::_V1::group", align 8
  %0 = load i64, ptr %_arg_dev_ptr1, align 8
  %1 = load i64, ptr %_arg_dev_ptr2, align 8
  %2 = load i64, ptr %_arg_dev_ptr3, align 8
  store i64 %2, ptr addrspace(3) @GKernel1, align 8
  store i64 %0, ptr addrspace(3) undef, align 8
  store i64 %1, ptr addrspace(3) undef, align 8
  %add.ptr.i = getelementptr inbounds i32, ptr addrspace(1) %_arg_dev_ptr, i64 %2
  store ptr addrspace(1) %add.ptr.i, ptr addrspace(3) undef, align 8
  %3 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalSize, align 32
  %4 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupSize, align 32
  %5 = load i64, ptr addrspace(1) @__spirv_BuiltInNumWorkgroups, align 32
  %6 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupId, align 32
  call void @llvm.lifetime.start.p0(i64 32, ptr nonnull %agg.tmp67)
  store i64 %3, ptr %agg.tmp67, align 1
  %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 8
  store i64 %4, ptr %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx, align 1
  %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 16
  store i64 %5, ptr %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx, align 1
  %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 24
  store i64 %6, ptr %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx, align 1
  %7 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationIndex, align 8
  tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  %cmpz27.i = icmp eq i64 %7, 0
  br i1 %cmpz27.i, label %leader.i, label %merge.i

leader.i:                                         ; preds = %entry
  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)
  br label %merge.i

merge.i:                                          ; preds = %leader.i, %entry
  tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  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)
  tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br i1 %cmpz27.i, label %wg_leader.i, label %wg_cf.i

wg_leader.i:                                      ; preds = %merge.i
  %g.ascast.i = addrspacecast ptr %agg.tmp67 to ptr addrspace(4)
  store ptr addrspace(4) %g.ascast.i, ptr addrspace(3) @GAsCast2, align 8
  store i32 0, ptr addrspace(3) @GCnt2, align 4
  br label %wg_cf.i

wg_cf.i:                                          ; preds = %wg_leader.i, %merge.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  %8 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalInvocationId, align 32
  %9 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationId, align 32
  %cmp.i.i.i.i.i.i = icmp ult i64 %8, 2147483648
  br label %for.cond.i

for.cond.i:                                       ; preds = %wg_cf18.i, %wg_cf.i
  %agg.tmp5.i.sroa.0.0 = phi i64 [ undef, %wg_cf.i ], [ %18, %wg_cf18.i ]
  %agg.tmp4.i.sroa.0.0 = phi i64 [ undef, %wg_cf.i ], [ %17, %wg_cf18.i ]
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br i1 %cmpz27.i, label %wg_leader8.i, label %wg_cf9.i

wg_leader8.i:                                     ; preds = %for.cond.i
  %10 = load i32, ptr addrspace(3) @GCnt2, align 4
  %cmp.i = icmp slt i32 %10, 2
  store i1 %cmp.i, ptr addrspace(3) @GCmp2, align 1
  br label %wg_cf9.i

wg_cf9.i:                                         ; preds = %wg_leader8.i, %for.cond.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  %wg_val_cmp.i = load i1, ptr addrspace(3) @GCmp2, align 1
  br i1 %wg_val_cmp.i, label %for.body.i, label %lexit2

for.body.i:                                       ; preds = %wg_cf9.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br i1 %cmpz27.i, label %TestMat25.i, label %LeaderMat22.i

TestMat25.i:                                      ; preds = %for.body.i
  store i64 %agg.tmp5.i.sroa.0.0, ptr addrspace(3) @WGCopy.6.0, align 8
  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
  store i64 5, ptr addrspace(3) @WGCopy.3.0, align 8
  store i64 %agg.tmp4.i.sroa.0.0, ptr addrspace(3) @WGCopy.5.0, align 8
  br label %LeaderMat22.i

LeaderMat22.i:                                    ; preds = %TestMat25.i, %for.body.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  %11 = load i64, ptr addrspace(3) @WGCopy.3.0, align 8
  %12 = load i64, ptr addrspace(3) @WGCopy.4.0, align 8
  %13 = inttoptr i64 %12 to ptr addrspace(4)
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef 2, i32 noundef 2, i32 noundef 272)
  %14 = getelementptr inbounds i8, ptr addrspace(4) %13, i64 24
  br label %for.cond.i.i

for.cond.i.i:                                     ; preds = %for.body.i.i, %LeaderMat22.i
  %storemerge.i.i = phi i64 [ %9, %LeaderMat22.i ], [ %add.i.i, %for.body.i.i ]
  %cmp.i.i = icmp ult i64 %storemerge.i.i, %11
  br i1 %cmp.i.i, label %for.body.i.i, label %lexit3

for.body.i.i:                                     ; preds = %for.cond.i.i
  call void @llvm.assume(i1 %cmp.i.i.i.i.i.i)
  %15 = load ptr addrspace(1), ptr addrspace(4) %14, align 8
  %arrayidx.i.i.i.i.i = getelementptr inbounds i32, ptr addrspace(1) %15, i64 %8
  %16 = load i32, ptr addrspace(1) %arrayidx.i.i.i.i.i, align 4
  %inc.i.i.i.i = add nsw i32 %16, 1
  store i32 %inc.i.i.i.i, ptr addrspace(1) %arrayidx.i.i.i.i.i, align 4
  %add.i.i = add i64 %storemerge.i.i, %4
  br label %for.cond.i.i

lexit3: ; preds = %for.cond.i.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef 2, i32 noundef 2, i32 noundef 272)
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br i1 %cmpz27.i, label %TestMat.i, label %LeaderMat.i

TestMat.i:                                        ; preds = %lexit3
  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
  store i64 %12, ptr addrspace(3) @WGCopy.4.0, align 8
  store i64 %11, ptr addrspace(3) @WGCopy.3.0, align 8
  store i64 2, ptr addrspace(3) @WGCopy.5.0, align 8
  br label %LeaderMat.i

LeaderMat.i:                                      ; preds = %TestMat.i, %lexit3
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  %17 = load i64, ptr addrspace(3) @WGCopy.5.0, align 8
  %18 = load i64, ptr addrspace(3) @WGCopy.6.0, align 8
  %19 = inttoptr i64 %18 to ptr addrspace(4)
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef 2, i32 noundef 2, i32 noundef 272)
  %20 = getelementptr inbounds i8, ptr addrspace(4) %19, i64 24
  br label %for.cond.i.i19

for.cond.i.i19:                                   ; preds = %for.body.i.i22, %LeaderMat.i
  %storemerge.i.i20 = phi i64 [ %9, %LeaderMat.i ], [ %add.i.i26, %for.body.i.i22 ]
  %cmp.i.i21 = icmp ult i64 %storemerge.i.i20, %17
  br i1 %cmp.i.i21, label %for.body.i.i22, label %lexit4

for.body.i.i22:                                   ; preds = %for.cond.i.i19
  call void @llvm.assume(i1 %cmp.i.i.i.i.i.i)
  %21 = load ptr addrspace(1), ptr addrspace(4) %20, align 8
  %arrayidx.i.i.i.i.i23 = getelementptr inbounds i32, ptr addrspace(1) %21, i64 %8
  %22 = load i32, ptr addrspace(1) %arrayidx.i.i.i.i.i23, align 4
  %inc.i.i.i.i25 = add nsw i32 %22, 1
  store i32 %inc.i.i.i.i25, ptr addrspace(1) %arrayidx.i.i.i.i.i23, align 4
  %add.i.i26 = add i64 %storemerge.i.i20, %4
  br label %for.cond.i.i19

lexit4: ; preds = %for.cond.i.i19
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef 2, i32 noundef 2, i32 noundef 272)
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br i1 %cmpz27.i, label %wg_leader17.i, label %wg_cf18.i

wg_leader17.i:                                    ; preds = %lexit4
  %23 = load i32, ptr addrspace(3) @GCnt2, align 4
  %inc.i = add nsw i32 %23, 1
  store i32 %inc.i, ptr addrspace(3) @GCnt2, align 4
  br label %wg_cf18.i

wg_cf18.i:                                        ; preds = %wg_leader17.i, %lexit4
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br label %for.cond.i

lexit2: ; preds = %wg_cf9.i
  call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %agg.tmp67)
  ret void
}

; Function Attrs: nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: write)
declare void @llvm.assume(i1 noundef)

; Function Attrs: convergent mustprogress norecurse nounwind
define 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) {
entry:
  %agg.tmp67 = alloca %"class.sycl::_V1::group", align 8
  %0 = load i64, ptr %_arg_dev_ptr1, align 8
  %1 = load i64, ptr %_arg_dev_ptr2, align 8
  %2 = load i64, ptr %_arg_dev_ptr3, align 8
  store i64 %2, ptr addrspace(3) @GKernel2, align 8
  store i64 %0, ptr addrspace(3) undef, align 8
  store i64 %1, ptr addrspace(3) undef, align 8
  %add.ptr.i = getelementptr inbounds i32, ptr addrspace(1) %_arg_dev_ptr, i64 %2
  store ptr addrspace(1) %add.ptr.i, ptr addrspace(3) undef, align 8
  %3 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalSize, align 32
  %4 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupSize, align 32
  %5 = load i64, ptr addrspace(1) @__spirv_BuiltInNumWorkgroups, align 32
  %6 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupId, align 32
  call void @llvm.lifetime.start.p0(i64 32, ptr nonnull %agg.tmp67)
  store i64 %3, ptr %agg.tmp67, align 1
  %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 8
  store i64 %4, ptr %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx, align 1
  %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 16
  store i64 %5, ptr %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx, align 1
  %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 24
  store i64 %6, ptr %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx, align 1
  %7 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationIndex, align 8
  tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  %cmpz15.i = icmp eq i64 %7, 0
  br i1 %cmpz15.i, label %leader.i, label %merge.i

leader.i:                                         ; preds = %entry
  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)
  br label %merge.i

merge.i:                                          ; preds = %leader.i, %entry
  tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  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)
  tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br i1 %cmpz15.i, label %wg_leader.i, label %wg_cf.i

wg_leader.i:                                      ; preds = %merge.i
  %g.ascast.i = addrspacecast ptr %agg.tmp67 to ptr addrspace(4)
  store ptr addrspace(4) %g.ascast.i, ptr addrspace(3) @GAscast3, align 8
  store i32 0, ptr addrspace(3) @GCnt3, align 4
  br label %wg_cf.i

wg_cf.i:                                          ; preds = %wg_leader.i, %merge.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  %8 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalInvocationId, align 32
  %9 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationId, align 32
  %cmp.i.i.i.i.i.i = icmp ult i64 %8, 2147483648
  br label %for.cond.i

for.cond.i:                                       ; preds = %wg_cf11.i, %wg_cf.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br i1 %cmpz15.i, label %wg_leader4.i, label %wg_cf5.i

wg_leader4.i:                                     ; preds = %for.cond.i
  %10 = load i32, ptr addrspace(3) @GCnt3, align 4
  %cmp.i = icmp slt i32 %10, 2
  store i1 %cmp.i, ptr addrspace(3) @GCmp3, align 1
  br label %wg_cf5.i

wg_cf5.i:                                         ; preds = %wg_leader4.i, %for.cond.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  %wg_val_cmp.i = load i1, ptr addrspace(3) @GCmp3, align 1
  br i1 %wg_val_cmp.i, label %for.body.i, label %lexit6

for.body.i:                                       ; preds = %wg_cf5.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br i1 %cmpz15.i, label %TestMat.i, label %LeaderMat.i

TestMat.i:                                        ; preds = %for.body.i
  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
  store i64 5, ptr addrspace(3) @WGCopy.9.0, align 8
  br label %LeaderMat.i

LeaderMat.i:                                      ; preds = %TestMat.i, %for.body.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  %11 = load i64, ptr addrspace(3) @WGCopy.9.0, align 8
  %12 = load i64, ptr addrspace(3) @WGCopy.10.0, align 8
  %13 = inttoptr i64 %12 to ptr addrspace(4)
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef 2, i32 noundef 2, i32 noundef 272)
  %14 = getelementptr inbounds i8, ptr addrspace(4) %13, i64 24
  br label %for.cond.i.i

for.cond.i.i:                                     ; preds = %for.body.i.i, %LeaderMat.i
  %storemerge.i.i = phi i64 [ %9, %LeaderMat.i ], [ %add.i.i, %for.body.i.i ]
  %cmp.i.i = icmp ult i64 %storemerge.i.i, %11
  br i1 %cmp.i.i, label %for.body.i.i, label %lexit7

for.body.i.i:                                     ; preds = %for.cond.i.i
  %cmp5.not.i.i.i.i.i.i = icmp ne i64 %storemerge.i.i, %9
  %cond.i.i.i.i = zext i1 %cmp5.not.i.i.i.i.i.i to i32
  call void @llvm.assume(i1 %cmp.i.i.i.i.i.i)
  %15 = load ptr addrspace(1), ptr addrspace(4) %14, align 8
  %arrayidx.i.i.i.i.i = getelementptr inbounds i32, ptr addrspace(1) %15, i64 %8
  %16 = load i32, ptr addrspace(1) %arrayidx.i.i.i.i.i, align 4
  %add.i.i.i.i = add nsw i32 %16, %cond.i.i.i.i
  store i32 %add.i.i.i.i, ptr addrspace(1) %arrayidx.i.i.i.i.i, align 4
  %add.i.i = add i64 %storemerge.i.i, %4
  br label %for.cond.i.i

lexit7: ; preds = %for.cond.i.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef 2, i32 noundef 2, i32 noundef 272)
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br i1 %cmpz15.i, label %wg_leader10.i, label %wg_cf11.i

wg_leader10.i:                                    ; preds = %lexit7
  %17 = load i32, ptr addrspace(3) @GCnt3, align 4
  %inc.i = add nsw i32 %17, 1
  store i32 %inc.i, ptr addrspace(3) @GCnt3, align 4
  br label %wg_cf11.i

wg_cf11.i:                                        ; preds = %wg_leader10.i, %lexit7
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br label %for.cond.i

lexit6: ; preds = %wg_cf5.i
  call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %agg.tmp67)
  ret void
}

; Function Attrs: convergent mustprogress norecurse nounwind
define 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) {
entry:
  %priv.i = alloca %"class.sycl::_V1::private_memory", align 4
  %agg.tmp67 = alloca %"class.sycl::_V1::group.15", align 8
  %0 = load i64, ptr %_arg_dev_ptr1, align 8
  %1 = load i64, ptr %_arg_dev_ptr2, align 8
  %2 = load i64, ptr %_arg_dev_ptr3, align 8
  store i64 %2, ptr addrspace(3) @GKernel3, align 8
  store i64 %0, ptr addrspace(3) undef, align 8
  store i64 %1, ptr addrspace(3) undef, align 8
  %add.ptr.i = getelementptr inbounds i32, ptr addrspace(1) %_arg_dev_ptr, i64 %2
  store ptr addrspace(1) %add.ptr.i, ptr addrspace(3) undef, align 8
  %3 = load i64, ptr addrspace(1) undef, align 8
  %4 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalSize, align 32
  %5 = load i64, ptr addrspace(1) undef, align 8
  %6 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupSize, align 32
  %7 = load i64, ptr addrspace(1) undef, align 8
  %8 = load i64, ptr addrspace(1) @__spirv_BuiltInNumWorkgroups, align 32
  %9 = load i64, ptr addrspace(1) undef, align 8
  %10 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupId, align 32
  call void @llvm.lifetime.start.p0(i64 64, ptr nonnull %agg.tmp67)
  store i64 %3, ptr %agg.tmp67, align 1
  %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 8
  store i64 %4, ptr %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx, align 1
  %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 16
  store i64 %5, ptr %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx, align 1
  %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 24
  store i64 %6, ptr %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx, align 1
  %agg.tmp6.sroa.5.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 32
  store i64 %7, ptr %agg.tmp6.sroa.5.0.agg.tmp67.sroa_idx, align 1
  %agg.tmp6.sroa.6.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 40
  store i64 %8, ptr %agg.tmp6.sroa.6.0.agg.tmp67.sroa_idx, align 1
  %agg.tmp6.sroa.7.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 48
  store i64 %9, ptr %agg.tmp6.sroa.7.0.agg.tmp67.sroa_idx, align 1
  %agg.tmp6.sroa.8.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 56
  store i64 %10, ptr %agg.tmp6.sroa.8.0.agg.tmp67.sroa_idx, align 1
  %11 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationIndex, align 8
  tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  %cmpz32.i = icmp eq i64 %11, 0
  br i1 %cmpz32.i, label %leader.i, label %merge.i

leader.i:                                         ; preds = %entry
  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)
  br label %merge.i

merge.i:                                          ; preds = %leader.i, %entry
  tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  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)
  %priv.ascast.i = addrspacecast ptr %priv.i to ptr addrspace(4)
  tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br i1 %cmpz32.i, label %wg_leader.i, label %wg_cf.i

wg_leader.i:                                      ; preds = %merge.i
  %g.ascast.i = addrspacecast ptr %agg.tmp67 to ptr addrspace(4)
  store ptr addrspace(4) %g.ascast.i, ptr addrspace(3) @GAsCast4, align 8
  call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %priv.i)
  store i32 0, ptr addrspace(3) @GCnt4, align 4
  br label %wg_cf.i

wg_cf.i:                                          ; preds = %wg_leader.i, %merge.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  %12 = load i64, ptr addrspace(1) undef, align 8
  %13 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalInvocationId, align 32
  %14 = load i64, ptr addrspace(1) undef, align 8
  %15 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationId, align 32
  %mul.i.i.i.i.i.i = mul i64 %12, %4
  %add.i.i.i.i.i.i = add i64 %mul.i.i.i.i.i.i, %13
  %cmp.i.i.i.i.i.i = icmp ult i64 %add.i.i.i.i.i.i, 2147483648
  %conv.i.i.i.i.i = trunc i64 %add.i.i.i.i.i.i to i32
  %y.i.i.i.i.i = getelementptr inbounds i8, ptr %priv.i, i64 4
  br label %for.cond.i

for.cond.i:                                       ; preds = %wg_cf20.i, %wg_cf.i
  %agg.tmp6.i.sroa.9.0 = phi ptr addrspace(4) [ undef, %wg_cf.i ], [ %agg.tmp6.i.sroa.9.0.copyload40, %wg_cf20.i ]
  %agg.tmp5.i.sroa.0.0 = phi i64 [ undef, %wg_cf.i ], [ %agg.tmp5.i.sroa.0.0.copyload44, %wg_cf20.i ]
  %agg.tmp5.i.sroa.8.0 = phi i64 [ undef, %wg_cf.i ], [ %agg.tmp5.i.sroa.8.0.copyload48, %wg_cf20.i ]
  %agg.tmp2.i.sroa.0.0 = phi ptr addrspace(4) [ undef, %wg_cf.i ], [ %agg.tmp2.i.sroa.0.0.copyload52, %wg_cf20.i ]
  %agg.tmp2.i.sroa.8.0 = phi ptr addrspace(4) [ undef, %wg_cf.i ], [ %agg.tmp2.i.sroa.8.0.copyload56, %wg_cf20.i ]
  %agg.tmp.i.sroa.0.0 = phi i64 [ undef, %wg_cf.i ], [ %agg.tmp.i.sroa.0.0.copyload60, %wg_cf20.i ]
  %agg.tmp.i.sroa.8.0 = phi i64 [ undef, %wg_cf.i ], [ %agg.tmp.i.sroa.8.0.copyload64, %wg_cf20.i ]
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br i1 %cmpz32.i, label %wg_leader10.i, label %wg_cf11.i

wg_leader10.i:                                    ; preds = %for.cond.i
  %16 = load i32, ptr addrspace(3) @GCnt4, align 4
  %cmp.i = icmp slt i32 %16, 2
  store i1 %cmp.i, ptr addrspace(3) @GCmp4, align 1
  br label %wg_cf11.i

wg_cf11.i:                                        ; preds = %wg_leader10.i, %for.cond.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  %wg_val_cmp.i = load i1, ptr addrspace(3) @GCmp4, align 1
  br i1 %wg_val_cmp.i, label %for.body.i, label %for.end.i

for.body.i:                                       ; preds = %wg_cf11.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br i1 %cmpz32.i, label %wg_leader13.i, label %wg_cf14.i

wg_leader13.i:                                    ; preds = %for.body.i
  br label %wg_cf14.i

wg_cf14.i:                                        ; preds = %wg_leader13.i, %for.body.i
  %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 ]
  %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 ]
  %agg.tmp.i.sroa.0.1 = phi i64 [ 7, %wg_leader13.i ], [ %agg.tmp.i.sroa.0.0, %for.body.i ]
  %agg.tmp.i.sroa.8.1 = phi i64 [ 3, %wg_leader13.i ], [ %agg.tmp.i.sroa.8.0, %for.body.i ]
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br i1 %cmpz32.i, label %TestMat30.i, label %LeaderMat27.i

TestMat30.i:                                      ; preds = %wg_cf14.i
  store i64 %agg.tmp.i.sroa.0.1, ptr addrspace(3) @WGCopy.13.0, align 8
  store i64 %agg.tmp.i.sroa.8.1, ptr addrspace(3) @WGCopy.13.1, align 8
  store ptr addrspace(4) %agg.tmp2.i.sroa.0.1, ptr addrspace(3) @WGCopy.14.0, align 8
  store ptr addrspace(4) %agg.tmp2.i.sroa.8.1, ptr addrspace(3) @WGCopy.14.1, align 8
  store i64 %agg.tmp5.i.sroa.0.0, ptr addrspace(3) @WGCopy.15.0, align 8
  store i64 %agg.tmp5.i.sroa.8.0, ptr addrspace(3) @WGCopy.15.1, align 8
  store ptr addrspace(4) %priv.ascast.i, ptr addrspace(3) @WGCopy.16.0, align 8
  store ptr addrspace(4) %agg.tmp6.i.sroa.9.0, ptr addrspace(3) @WGCopy.16.1, align 8
  br label %LeaderMat27.i

LeaderMat27.i:                                    ; preds = %TestMat30.i, %wg_cf14.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  %agg.tmp6.i.sroa.0.0.copyload = load ptr addrspace(4), ptr addrspace(3) @WGCopy.16.0, align 8
  %agg.tmp6.i.sroa.9.0.copyload = load ptr addrspace(4), ptr addrspace(3) @WGCopy.16.1, align 8
  %agg.tmp5.i.sroa.0.0.copyload = load i64, ptr addrspace(3) @WGCopy.15.0, align 8
  %agg.tmp5.i.sroa.8.0.copyload = load i64, ptr addrspace(3) @WGCopy.15.1, align 8
  %agg.tmp2.i.sroa.0.0.copyload = load ptr addrspace(4), ptr addrspace(3) @WGCopy.14.0, align 8
  %agg.tmp.i.sroa.0.0.copyload = load i64, ptr addrspace(3) @WGCopy.13.0, align 8
  %agg.tmp.i.sroa.8.0.copyload = load i64, ptr addrspace(3) @WGCopy.13.1, align 8
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef 2, i32 noundef 2, i32 noundef 272)
  %17 = getelementptr inbounds i8, ptr addrspace(4) %agg.tmp2.i.sroa.0.0.copyload, i64 24
  br label %for.cond.i.i

for.cond.i.i:                                     ; preds = %lexit10, %LeaderMat27.i
  %storemerge.i.i = phi i64 [ %14, %LeaderMat27.i ], [ %add.i.i, %lexit10 ]
  %cmp.i.i = icmp ult i64 %storemerge.i.i, %agg.tmp.i.sroa.0.0.copyload
  br i1 %cmp.i.i, label %for.cond.i.i.i, label %lexit11

for.cond.i.i.i:                                   ; preds = %for.body.i.i.i, %for.cond.i.i
  %storemerge.i.i.i = phi i64 [ %add.i.i.i, %for.body.i.i.i ], [ %15, %for.cond.i.i ]
  %cmp.i.i.i = icmp ult i64 %storemerge.i.i.i, %agg.tmp.i.sroa.8.0.copyload
  br i1 %cmp.i.i.i, label %for.body.i.i.i, label %lexit10

for.body.i.i.i:                                   ; preds = %for.cond.i.i.i
  call void @llvm.assume(i1 %cmp.i.i.i.i.i.i)
  %18 = load ptr addrspace(1), ptr addrspace(4) %17, align 8
  %arrayidx.i.i.i.i.i.i = getelementptr inbounds i32, ptr addrspace(1) %18, i64 %add.i.i.i.i.i.i
  %19 = load i32, ptr addrspace(1) %arrayidx.i.i.i.i.i.i, align 4
  %inc.i.i.i.i.i = add nsw i32 %19, 1
  store i32 %inc.i.i.i.i.i, ptr addrspace(1) %arrayidx.i.i.i.i.i.i, align 4
  store i32 %conv.i.i.i.i.i, ptr %priv.i, align 4
  store i32 5, ptr %y.i.i.i.i.i, align 4
  %add.i.i.i = add i64 %storemerge.i.i.i, %6
  br label %for.cond.i.i.i

lexit10: ; preds = %for.cond.i.i.i
  %add.i.i = add i64 %storemerge.i.i, %5
  br label %for.cond.i.i

lexit11: ; preds = %for.cond.i.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef 2, i32 noundef 2, i32 noundef 272)
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br i1 %cmpz32.i, label %wg_leader16.i, label %wg_cf17.i

wg_leader16.i:                                    ; preds = %lexit11
  br label %wg_cf17.i

wg_cf17.i:                                        ; preds = %wg_leader16.i, %lexit11
  %agg.tmp6.i.sroa.0.1 = phi ptr addrspace(4) [ %priv.ascast.i, %wg_leader16.i ], [ %agg.tmp6.i.sroa.0.0.copyload, %lexit11 ]
  %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 ]
  %agg.tmp5.i.sroa.0.1 = phi i64 [ 7, %wg_leader16.i ], [ %agg.tmp5.i.sroa.0.0.copyload, %lexit11 ]
  %agg.tmp5.i.sroa.8.1 = phi i64 [ 3, %wg_leader16.i ], [ %agg.tmp5.i.sroa.8.0.copyload, %lexit11 ]
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br i1 %cmpz32.i, label %TestMat.i, label %LeaderMat.i

TestMat.i:                                        ; preds = %wg_cf17.i
  store i64 %agg.tmp.i.sroa.0.0.copyload, ptr addrspace(3) @WGCopy.13.0, align 8
  store i64 %agg.tmp.i.sroa.8.0.copyload, ptr addrspace(3) @WGCopy.13.1, align 8
  store ptr addrspace(4) %agg.tmp2.i.sroa.0.0.copyload, ptr addrspace(3) @WGCopy.14.0, align 8
  store ptr addrspace(4) %priv.ascast.i, ptr addrspace(3) @WGCopy.14.1, align 8
  store i64 %agg.tmp5.i.sroa.0.1, ptr addrspace(3) @WGCopy.15.0, align 8
  store i64 %agg.tmp5.i.sroa.8.1, ptr addrspace(3) @WGCopy.15.1, align 8
  store ptr addrspace(4) %agg.tmp6.i.sroa.0.1, ptr addrspace(3) @WGCopy.16.0, align 8
  store ptr addrspace(4) %agg.tmp6.i.sroa.9.1, ptr addrspace(3) @WGCopy.16.1, align 8
  br label %LeaderMat.i

LeaderMat.i:                                      ; preds = %TestMat.i, %wg_cf17.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  %agg.tmp6.i.sroa.9.0.copyload40 = load ptr addrspace(4), ptr addrspace(3) @WGCopy.16.1, align 8
  %agg.tmp5.i.sroa.0.0.copyload44 = load i64, ptr addrspace(3) @WGCopy.15.0, align 8
  %agg.tmp5.i.sroa.8.0.copyload48 = load i64, ptr addrspace(3) @WGCopy.15.1, align 8
  %agg.tmp2.i.sroa.0.0.copyload52 = load ptr addrspace(4), ptr addrspace(3) @WGCopy.14.0, align 8
  %agg.tmp2.i.sroa.8.0.copyload56 = load ptr addrspace(4), ptr addrspace(3) @WGCopy.14.1, align 8
  %agg.tmp.i.sroa.0.0.copyload60 = load i64, ptr addrspace(3) @WGCopy.13.0, align 8
  %agg.tmp.i.sroa.8.0.copyload64 = load i64, ptr addrspace(3) @WGCopy.13.1, align 8
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef 2, i32 noundef 2, i32 noundef 272)
  %20 = getelementptr inbounds i8, ptr addrspace(4) %agg.tmp6.i.sroa.9.0.copyload40, i64 24
  br label %for.cond.i.i25

for.cond.i.i25:                                   ; preds = %lexit12, %LeaderMat.i
  %storemerge.i.i26 = phi i64 [ %14, %LeaderMat.i ], [ %add.i.i31, %lexit12 ]
  %cmp.i.i27 = icmp ult i64 %storemerge.i.i26, %agg.tmp5.i.sroa.0.0.copyload44
  br i1 %cmp.i.i27, label %for.cond.i.i.i28, label %lexit13

for.cond.i.i.i28:                                 ; preds = %for.body.i.i.i32, %for.cond.i.i25
  %storemerge.i.i.i29 = phi i64 [ %add.i.i.i35, %for.body.i.i.i32 ], [ %15, %for.cond.i.i25 ]
  %cmp.i.i.i30 = icmp ult i64 %storemerge.i.i.i29, %agg.tmp5.i.sroa.8.0.copyload48
  br i1 %cmp.i.i.i30, label %for.body.i.i.i32, label %lexit12

for.body.i.i.i32:                                 ; preds = %for.cond.i.i.i28
  %21 = load i32, ptr %priv.i, align 4
  %22 = load i32, ptr %y.i.i.i.i.i, align 4
  %add.i.i.i.i.i = add nsw i32 %21, %22
  call void @llvm.assume(i1 %cmp.i.i.i.i.i.i)
  %23 = load ptr addrspace(1), ptr addrspace(4) %20, align 8
  %arrayidx.i.i.i.i.i.i33 = getelementptr inbounds i32, ptr addrspace(1) %23, i64 %add.i.i.i.i.i.i
  %24 = load i32, ptr addrspace(1) %arrayidx.i.i.i.i.i.i33, align 4
  %add4.i.i.i.i.i = add nsw i32 %24, %add.i.i.i.i.i
  store i32 %add4.i.i.i.i.i, ptr addrspace(1) %arrayidx.i.i.i.i.i.i33, align 4
  %add.i.i.i35 = add i64 %storemerge.i.i.i29, %6
  br label %for.cond.i.i.i28

lexit12: ; preds = %for.cond.i.i.i28
  %add.i.i31 = add i64 %storemerge.i.i26, %5
  br label %for.cond.i.i25

lexit13: ; preds = %for.cond.i.i25
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef 2, i32 noundef 2, i32 noundef 272)
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br i1 %cmpz32.i, label %wg_leader19.i, label %wg_cf20.i

wg_leader19.i:                                    ; preds = %lexit13
  %25 = load i32, ptr addrspace(3) @GCnt4, align 4
  %inc.i = add nsw i32 %25, 1
  store i32 %inc.i, ptr addrspace(3) @GCnt4, align 4
  br label %wg_cf20.i

wg_cf20.i:                                        ; preds = %wg_leader19.i, %lexit13
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br label %for.cond.i

for.end.i:                                        ; preds = %wg_cf11.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br i1 %cmpz32.i, label %wg_leader22.i, label %lexit14

wg_leader22.i:                                    ; preds = %for.end.i
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %priv.i)
  br label %lexit14

lexit14: ; preds = %wg_leader22.i, %for.end.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  call void @llvm.lifetime.end.p0(i64 64, ptr nonnull %agg.tmp67)
  ret void
}

; Function Attrs: convergent mustprogress norecurse nounwind
define 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) {
entry:
  %agg.tmp67 = alloca %"class.sycl::_V1::group", align 8
  %0 = load i64, ptr %_arg_dev_ptr1, align 8
  %1 = load i64, ptr %_arg_dev_ptr2, align 8
  %2 = load i64, ptr %_arg_dev_ptr3, align 8
  store i64 %2, ptr addrspace(3) @GKernel4, align 8
  store i64 %0, ptr addrspace(3) undef, align 8
  store i64 %1, ptr addrspace(3) undef, align 8
  %add.ptr.i = getelementptr inbounds i32, ptr addrspace(1) %_arg_dev_ptr, i64 %2
  store ptr addrspace(1) %add.ptr.i, ptr addrspace(3) undef, align 8
  %3 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalSize, align 32
  %4 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupSize, align 32
  %5 = load i64, ptr addrspace(1) @__spirv_BuiltInNumWorkgroups, align 32
  %6 = load i64, ptr addrspace(1) @__spirv_BuiltInWorkgroupId, align 32
  call void @llvm.lifetime.start.p0(i64 32, ptr nonnull %agg.tmp67)
  store i64 %3, ptr %agg.tmp67, align 1
  %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 8
  store i64 %4, ptr %agg.tmp6.sroa.2.0.agg.tmp67.sroa_idx, align 1
  %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 16
  store i64 %5, ptr %agg.tmp6.sroa.3.0.agg.tmp67.sroa_idx, align 1
  %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx = getelementptr inbounds i8, ptr %agg.tmp67, i64 24
  store i64 %6, ptr %agg.tmp6.sroa.4.0.agg.tmp67.sroa_idx, align 1
  %7 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationIndex, align 8
  tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  %cmpz16.i = icmp eq i64 %7, 0
  br i1 %cmpz16.i, label %leader.i, label %merge.i

leader.i:                                         ; preds = %entry
  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)
  br label %merge.i

merge.i:                                          ; preds = %leader.i, %entry
  tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  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)
  tail call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br i1 %cmpz16.i, label %wg_leader.i, label %wg_cf.i

wg_leader.i:                                      ; preds = %merge.i
  %g.ascast.i = addrspacecast ptr %agg.tmp67 to ptr addrspace(4)
  store ptr addrspace(4) %g.ascast.i, ptr addrspace(3) @GAsCast5, align 8
  store i32 0, ptr addrspace(3) @GCnt5, align 4
  br label %wg_cf.i

wg_cf.i:                                          ; preds = %wg_leader.i, %merge.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  %wg_val_g.ascast.i = load ptr addrspace(4), ptr addrspace(3) @GAsCast5, align 8
  %8 = load i64, ptr addrspace(1) @__spirv_BuiltInLocalInvocationId, align 32
  %9 = trunc i64 %4 to i32
  br label %for.cond.i

for.cond.i:                                       ; preds = %wg_cf12.i, %wg_cf.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br i1 %cmpz16.i, label %wg_leader5.i, label %wg_cf6.i

wg_leader5.i:                                     ; preds = %for.cond.i
  %10 = load i32, ptr addrspace(3) @GCnt5, align 4
  %cmp.i = icmp slt i32 %10, 2
  store i1 %cmp.i, ptr addrspace(3) @GCmp5, align 1
  br label %wg_cf6.i

wg_cf6.i:                                         ; preds = %wg_leader5.i, %for.cond.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  %wg_val_cmp.i = load i1, ptr addrspace(3) @GCmp5, align 1
  br i1 %wg_val_cmp.i, label %for.body.i, label %lexit20

for.body.i:                                       ; preds = %wg_cf6.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br i1 %cmpz16.i, label %TestMat.i, label %LeaderMat.i

TestMat.i:                                        ; preds = %for.body.i
  store ptr addrspace(4) %wg_val_g.ascast.i, ptr addrspace(3) @WGCopy.20.0, align 8
  store ptr addrspace(4) addrspacecast (ptr addrspace(3) @GKernel4 to ptr addrspace(4)), ptr addrspace(3) @WGCopy.20.1, align 8
  store i64 5, ptr addrspace(3) @WGCopy.19.0, align 8
  br label %LeaderMat.i

LeaderMat.i:                                      ; preds = %TestMat.i, %for.body.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  %11 = load i64, ptr addrspace(3) @WGCopy.19.0, align 8
  %agg.tmp2.i.sroa.0.0.copyload = load ptr addrspace(4), ptr addrspace(3) @WGCopy.20.0, align 8
  %agg.tmp2.i.sroa.6.0.copyload = load ptr addrspace(4), ptr addrspace(3) @WGCopy.20.1, align 8
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef 2, i32 noundef 2, i32 noundef 272)
  %index.i.i.i.i.i = getelementptr inbounds i8, ptr addrspace(4) %agg.tmp2.i.sroa.0.0.copyload, i64 24
  %12 = getelementptr inbounds i8, ptr addrspace(4) %agg.tmp2.i.sroa.6.0.copyload, i64 24
  %13 = trunc i64 %11 to i32
  br label %for.cond.i.i

for.cond.i.i:                                     ; preds = %for.body.i.i, %LeaderMat.i
  %storemerge.i.i = phi i64 [ %8, %LeaderMat.i ], [ %add.i.i, %for.body.i.i ]
  %cmp.i.i = icmp ult i64 %storemerge.i.i, %11
  br i1 %cmp.i.i, label %for.body.i.i, label %lexit21

for.body.i.i:                                     ; preds = %for.cond.i.i
  %14 = load i64, ptr addrspace(4) %index.i.i.i.i.i, align 8
  %mul.i.i.i.i = mul i64 %14, 10
  %mul3.i.i.i.i = shl i64 %storemerge.i.i, 1
  %add.i.i.i.i = add i64 %mul.i.i.i.i, %mul3.i.i.i.i
  %15 = load ptr addrspace(1), ptr addrspace(4) %12, align 8
  %arrayidx.i.i.i.i.i = getelementptr inbounds i32, ptr addrspace(1) %15, i64 %add.i.i.i.i
  %16 = load i32, ptr addrspace(1) %arrayidx.i.i.i.i.i, align 4
  %conv9.i.i.i.i = add i32 %16, %13
  store i32 %conv9.i.i.i.i, ptr addrspace(1) %arrayidx.i.i.i.i.i, align 4
  %add14.i.i.i.i = or disjoint i64 %add.i.i.i.i, 1
  %17 = load ptr addrspace(1), ptr addrspace(4) %12, align 8
  %arrayidx.i25.i.i.i.i = getelementptr inbounds i32, ptr addrspace(1) %17, i64 %add14.i.i.i.i
  %18 = load i32, ptr addrspace(1) %arrayidx.i25.i.i.i.i, align 4
  %conv18.i.i.i.i = add i32 %18, %9
  store i32 %conv18.i.i.i.i, ptr addrspace(1) %arrayidx.i25.i.i.i.i, align 4
  %add.i.i = add i64 %storemerge.i.i, %4
  br label %for.cond.i.i

lexit21: ; preds = %for.cond.i.i
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 noundef 2, i32 noundef 2, i32 noundef 272)
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br i1 %cmpz16.i, label %wg_leader11.i, label %wg_cf12.i

wg_leader11.i:                                    ; preds = %lexit21
  %19 = load i32, ptr addrspace(3) @GCnt5, align 4
  %inc.i = add nsw i32 %19, 1
  store i32 %inc.i, ptr addrspace(3) @GCnt5, align 4
  br label %wg_cf12.i

wg_cf12.i:                                        ; preds = %wg_leader11.i, %lexit21
  call spir_func void @_Z22__spirv_ControlBarrierjjj(i32 2, i32 2, i32 272)
  br label %for.cond.i

lexit20: ; preds = %wg_cf6.i
  call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %agg.tmp67)
  ret void
}