benchmarks-bfs-dump.promote.ll

; ModuleID = '<stdin>'
target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
target triple = "x86_64-unknown-linux-gnu"

%struct.LocalQueues = type { [8 x i32], [8 x [400 x i32]], [8 x i32] }
%"class.hc::short_vector::int_2.0" = type { i32, i32 }
%struct.grid_launch_parm = type { %struct.gl_dim3, %struct.gl_dim3, %struct.gl_dim3, %struct.gl_dim3, i32, %"class.hc::accelerator_view"*, %"class.hc::completion_future"* }
%struct.gl_dim3 = type { i32, i32, i32 }
%"class.hc::accelerator_view" = type { %"class.std::__1::shared_ptr" }
%"class.std::__1::shared_ptr" = type { %"class.Kalmar::KalmarQueue"*, %"class.std::__1::__shared_weak_count"* }
%"class.Kalmar::KalmarQueue" = type { i32 (...)**, %"class.Kalmar::KalmarDevice"*, i32, i32 }
%"class.Kalmar::KalmarDevice" = type { i32 (...)**, i32, %"class.std::__1::map", %"class.std::__1::mutex" }
%"class.std::__1::map" = type { %"class.std::__1::__tree.13" }
%"class.std::__1::__tree.13" = type { %"class.std::__1::__tree_node.14"*, %"class.std::__1::__compressed_pair", %"class.std::__1::__compressed_pair.20" }
%"class.std::__1::__tree_node.14" = type { %"class.std::__1::__tree_node_base.base", %"union.std::__1::__value_type" }
%"class.std::__1::__tree_node_base.base" = type <{ %"class.std::__1::__tree_end_node", %"class.std::__1::__tree_node_base"*, %"class.std::__1::__tree_node_base"*, i8 }>
%"class.std::__1::__tree_end_node" = type { %"class.std::__1::__tree_node_base"* }
%"class.std::__1::__tree_node_base" = type { %"class.std::__1::__tree_end_node", %"class.std::__1::__tree_node_base"*, %"class.std::__1::__tree_node_base"*, i8 }
%"union.std::__1::__value_type" = type { %"struct.std::__1::pair" }
%"struct.std::__1::pair" = type { %"class.std::__1::__thread_id", %"class.std::__1::shared_ptr" }
%"class.std::__1::__thread_id" = type { i64 }
%"class.std::__1::__compressed_pair" = type { %"class.std::__1::__libcpp_compressed_pair_imp" }
%"class.std::__1::__libcpp_compressed_pair_imp" = type { %"class.std::__1::__tree_end_node" }
%"class.std::__1::__compressed_pair.20" = type { %"class.std::__1::__libcpp_compressed_pair_imp.21" }
%"class.std::__1::__libcpp_compressed_pair_imp.21" = type { i64 }
%"class.std::__1::mutex" = type { %union.pthread_mutex_t }
%union.pthread_mutex_t = type { %"struct.(anonymous union)::__pthread_mutex_s" }
%"struct.(anonymous union)::__pthread_mutex_s" = type { i32, i32, i32, i32, i32, i16, i16, %struct.__pthread_internal_list }
%struct.__pthread_internal_list = type { %struct.__pthread_internal_list*, %struct.__pthread_internal_list* }
%"class.std::__1::__shared_weak_count" = type { %"class.std::__1::__shared_count", i64 }
%"class.std::__1::__shared_count" = type { i32 (...)**, i64 }
%"class.hc::completion_future" = type { %"class.std::__1::shared_future", %"class.std::__1::thread"*, %"class.std::__1::shared_ptr.24" }
%"class.std::__1::shared_future" = type { %"class.std::__1::__assoc_sub_state"* }
%"class.std::__1::__assoc_sub_state" = type { %"class.std::__1::__shared_count", %"class.std::exception_ptr", %"class.std::__1::mutex", %"class.std::__1::condition_variable", i32 }
%"class.std::exception_ptr" = type { i8* }
%"class.std::__1::condition_variable" = type { %union.pthread_cond_t }
%union.pthread_cond_t = type { %struct.anon }
%struct.anon = type { i32, i32, i64, i64, i64, i8*, i32, i32 }
%"class.std::__1::thread" = type { i64 }
%"class.std::__1::shared_ptr.24" = type { %"class.Kalmar::KalmarAsyncOp"*, %"class.std::__1::__shared_weak_count"* }
%"class.Kalmar::KalmarAsyncOp" = type { i32 (...)** }

@count24 = addrspace(1) global i32 0, align 4
@no_of_nodes_vol27 = addrspace(1) global i32 0, align 4
@stay_vol29 = addrspace(1) global i32 0, align 4
@ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q245 = internal addrspace(3) global %struct.LocalQueues undef, section "clamp_opencl_local", align 4
@ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E8prefix_q346 = internal addrspace(3) global [8 x i32] undef, section "clamp_opencl_local", align 16
@ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7next_wf447 = internal addrspace(3) global [512 x i32] undef, section "clamp_opencl_local", align 16
@ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7tot_sum548 = internal unnamed_addr addrspace(3) global i32 undef, section "clamp_opencl_local", align 4
@ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E7local_q631 = internal addrspace(3) global %struct.LocalQueues undef, section "clamp_opencl_local", align 4
@ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E8prefix_q732 = internal addrspace(3) global [8 x i32] undef, section "clamp_opencl_local", align 16
@ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E5shift833 = internal unnamed_addr addrspace(3) global i32 undef, section "clamp_opencl_local", align 4
@ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E14no_of_nodes_sm934 = internal unnamed_addr addrspace(3) global i32 undef, section "clamp_opencl_local", align 4
@ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E8odd_time1035 = internal unnamed_addr addrspace(3) global i32 undef, section "clamp_opencl_local", align 4
@ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q1116 = internal addrspace(3) global %struct.LocalQueues undef, section "clamp_opencl_local", align 4
@ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E8prefix_q1217 = internal addrspace(3) global [8 x i32] undef, section "clamp_opencl_local", align 16
@ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E5shift1318 = internal unnamed_addr addrspace(3) global i32 undef, section "clamp_opencl_local", align 4

; Function Attrs: nounwind readnone
declare i64 @amp_get_global_id(i32) #0

; Function Attrs: nounwind
declare void @llvm.memcpy.p0i8.p0i8.i64(i8* nocapture, i8* nocapture readonly, i64, i32, i1) #1

; Function Attrs: nounwind
declare void @llvm.lifetime.start(i64, i8* nocapture) #1

; Function Attrs: nounwind
declare void @llvm.lifetime.end(i64, i8* nocapture) #1

; Function Attrs: nounwind readnone
declare i64 @amp_get_local_size(i32) #0

; Function Attrs: nounwind readnone
declare i64 @hc_get_num_groups(i32) #0

; Function Attrs: noduplicate
declare void @hc_barrier(i32) #2

; Function Attrs: nounwind readnone
declare i64 @amp_get_local_id(i32) #3

; Function Attrs: nounwind readnone
declare i64 @hc_get_group_id(i32) #3

; Function Attrs: nounwind readnone
declare i64 @amp_get_group_id(i32) #0

; Function Attrs: nounwind uwtable
define spir_kernel void @ZZ16ihipMemsetKernelIcEN2hc17completion_futureEP12ihipStream_tPT_S4_mENUlNS0_11tiled_indexILi1EEEE_19__cxxamp_trampolineEmPcc(i64, i8 addrspace(1)*, i8 signext) #4 align 2 {
  %4 = tail call spir_func i64 @amp_get_global_id(i32 0) #7
  %sext.i = shl i64 %4, 32
  %5 = ashr exact i64 %sext.i, 32
  %6 = icmp ult i64 %5, %0
  br i1 %6, label %.lr.ph.i, label %_ZZ16ihipMemsetKernelIcEN2hc17completion_futureEP12ihipStream_tPT_S4_mENKUlNS0_11tiled_indexILi1EEEE_clES7_.exit

.lr.ph.i:                                         ; preds = %3
  %7 = tail call spir_func i64 @amp_get_local_size(i32 0) #7
  %8 = tail call spir_func i64 @hc_get_num_groups(i32 0) #7
  %9 = shl i64 %7, 32
  %sext3.i = mul i64 %9, %8
  %10 = ashr exact i64 %sext3.i, 32
  br label %11

; <label>:11                                      ; preds = %11, %.lr.ph.i
  %indvars.iv.i = phi i64 [ %5, %.lr.ph.i ], [ %indvars.iv.next.i, %11 ]
  %12 = getelementptr inbounds i8 addrspace(1)* %1, i64 %indvars.iv.i
  store i8 %2, i8 addrspace(1)* %12, align 1, !tbaa !23
  %indvars.iv.next.i = add nsw i64 %indvars.iv.i, %10
  %13 = icmp ult i64 %indvars.iv.next.i, %0
  br i1 %13, label %11, label %_ZZ16ihipMemsetKernelIcEN2hc17completion_futureEP12ihipStream_tPT_S4_mENKUlNS0_11tiled_indexILi1EEEE_clES7_.exit.loopexit

_ZZ16ihipMemsetKernelIcEN2hc17completion_futureEP12ihipStream_tPT_S4_mENKUlNS0_11tiled_indexILi1EEEE_clES7_.exit.loopexit: ; preds = %11
  br label %_ZZ16ihipMemsetKernelIcEN2hc17completion_futureEP12ihipStream_tPT_S4_mENKUlNS0_11tiled_indexILi1EEEE_clES7_.exit

_ZZ16ihipMemsetKernelIcEN2hc17completion_futureEP12ihipStream_tPT_S4_mENKUlNS0_11tiled_indexILi1EEEE_clES7_.exit: ; preds = %_ZZ16ihipMemsetKernelIcEN2hc17completion_futureEP12ihipStream_tPT_S4_mENKUlNS0_11tiled_indexILi1EEEE_clES7_.exit.loopexit, %3
  ret void
}

; Function Attrs: nounwind uwtable
define spir_kernel void @ZZ16ihipMemsetKernelIjEN2hc17completion_futureEP12ihipStream_tPT_S4_mENUlNS0_11tiled_indexILi1EEEE_19__cxxamp_trampolineEmPjj(i64, i32 addrspace(1)*, i32) #4 align 2 {
  %4 = tail call spir_func i64 @amp_get_global_id(i32 0) #7
  %sext.i = shl i64 %4, 32
  %5 = ashr exact i64 %sext.i, 32
  %6 = icmp ult i64 %5, %0
  br i1 %6, label %.lr.ph.i, label %_ZZ16ihipMemsetKernelIjEN2hc17completion_futureEP12ihipStream_tPT_S4_mENKUlNS0_11tiled_indexILi1EEEE_clES7_.exit

.lr.ph.i:                                         ; preds = %3
  %7 = tail call spir_func i64 @amp_get_local_size(i32 0) #7
  %8 = tail call spir_func i64 @hc_get_num_groups(i32 0) #7
  %9 = shl i64 %7, 32
  %sext3.i = mul i64 %9, %8
  %10 = ashr exact i64 %sext3.i, 32
  br label %11

; <label>:11                                      ; preds = %11, %.lr.ph.i
  %indvars.iv.i = phi i64 [ %5, %.lr.ph.i ], [ %indvars.iv.next.i, %11 ]
  %12 = getelementptr inbounds i32 addrspace(1)* %1, i64 %indvars.iv.i
  store i32 %2, i32 addrspace(1)* %12, align 4, !tbaa !26
  %indvars.iv.next.i = add nsw i64 %indvars.iv.i, %10
  %13 = icmp ult i64 %indvars.iv.next.i, %0
  br i1 %13, label %11, label %_ZZ16ihipMemsetKernelIjEN2hc17completion_futureEP12ihipStream_tPT_S4_mENKUlNS0_11tiled_indexILi1EEEE_clES7_.exit.loopexit

_ZZ16ihipMemsetKernelIjEN2hc17completion_futureEP12ihipStream_tPT_S4_mENKUlNS0_11tiled_indexILi1EEEE_clES7_.exit.loopexit: ; preds = %11
  br label %_ZZ16ihipMemsetKernelIjEN2hc17completion_futureEP12ihipStream_tPT_S4_mENKUlNS0_11tiled_indexILi1EEEE_clES7_.exit

_ZZ16ihipMemsetKernelIjEN2hc17completion_futureEP12ihipStream_tPT_S4_mENKUlNS0_11tiled_indexILi1EEEE_clES7_.exit: ; preds = %_ZZ16ihipMemsetKernelIjEN2hc17completion_futureEP12ihipStream_tPT_S4_mENKUlNS0_11tiled_indexILi1EEEE_clES7_.exit.loopexit, %3
  ret void
}

; Function Attrs: uwtable
define spir_kernel void @ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_(i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32 addrspace(1)*, i32 addrspace(1)*, %"class.hc::short_vector::int_2.0" addrspace(1)*, %"class.hc::short_vector::int_2.0" addrspace(1)*, i32 addrspace(1)*, i32 addrspace(1)*, i32, i32 addrspace(1)*, i32, i32, i32 addrspace(1)*) #5 align 2 {
  %25 = alloca %struct.grid_launch_parm, align 8
  %26 = alloca %struct.grid_launch_parm, align 16
  %.sroa.23 = alloca [20 x i8], align 4
  %27 = insertelement <4 x i32> undef, i32 %0, i32 0
  %28 = insertelement <4 x i32> %27, i32 %1, i32 1
  %29 = insertelement <4 x i32> %28, i32 %2, i32 2
  %30 = insertelement <4 x i32> %29, i32 %3, i32 3
  %31 = insertelement <4 x i32> undef, i32 %4, i32 0
  %32 = insertelement <4 x i32> %31, i32 %5, i32 1
  %33 = tail call spir_func i64 @amp_get_local_id(i32 2) #7
  %34 = tail call spir_func i64 @amp_get_local_id(i32 1) #7
  %35 = tail call spir_func i64 @amp_get_local_id(i32 0) #7
  %36 = trunc i64 %33 to i32
  %37 = trunc i64 %34 to i32
  %38 = trunc i64 %35 to i32
  %39 = tail call spir_func i64 @amp_get_group_id(i32 2) #7
  %40 = tail call spir_func i64 @amp_get_group_id(i32 1) #7
  %41 = tail call spir_func i64 @amp_get_group_id(i32 0) #7
  %42 = trunc i64 %39 to i32
  %43 = trunc i64 %40 to i32
  %44 = trunc i64 %41 to i32
  %45 = bitcast %struct.grid_launch_parm* %26 to i8*
  call spir_func void @llvm.lifetime.start(i64 72, i8* %45)
  %46 = insertelement <4 x i32> %32, i32 %44, i32 2
  %47 = insertelement <4 x i32> %46, i32 %43, i32 3
  %48 = insertelement <4 x i32> undef, i32 %42, i32 0
  %49 = insertelement <4 x i32> %48, i32 %38, i32 1
  %50 = insertelement <4 x i32> %49, i32 %37, i32 2
  %51 = insertelement <4 x i32> %50, i32 %36, i32 3
  %52 = bitcast %struct.grid_launch_parm* %26 to <4 x i32>*
  store <4 x i32> %30, <4 x i32>* %52, align 16
  %53 = getelementptr inbounds %struct.grid_launch_parm* %26, i64 0, i32 1, i32 1
  %54 = bitcast i32* %53 to <4 x i32>*
  store <4 x i32> %47, <4 x i32>* %54, align 8
  %55 = getelementptr inbounds %struct.grid_launch_parm* %26, i64 0, i32 2, i32 2
  %56 = bitcast i32* %55 to <4 x i32>*
  store <4 x i32> %51, <4 x i32>* %56, align 8
  %57 = getelementptr inbounds %struct.grid_launch_parm* %26, i64 0, i32 4
  store i32 %12, i32* %57, align 16
  %58 = getelementptr inbounds i8* %45, i64 52
  %59 = getelementptr inbounds [20 x i8]* %.sroa.23, i64 0, i64 0
  call spir_func void @llvm.memcpy.p0i8.p0i8.i64(i8* %58, i8* %59, i64 20, i32 4, i1 false)
  %60 = bitcast %struct.grid_launch_parm* %25 to i8*
  call spir_func void @llvm.lifetime.start(i64 72, i8* %60)
  %tmp = bitcast %struct.grid_launch_parm* %25 to i8*
  %tmp1 = bitcast %struct.grid_launch_parm* %26 to i8*
  call spir_func void @llvm.memcpy.p0i8.p0i8.i64(i8* %tmp, i8* %tmp1, i64 72, i32 1, i1 false)
  %61 = call spir_func i64 @amp_get_local_id(i32 0) #7
  %62 = icmp slt i64 %61, 8
  br i1 %62, label %63, label %74

; <label>:63                                      ; preds = %24
  %sext2.i = shl i64 %61, 32
  %64 = ashr exact i64 %sext2.i, 32
  %65 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q1116, i64 0, i32 0, i64 %64
  store i32 0, i32 addrspace(3)* %65, align 4, !tbaa !26
  %66 = call spir_func i64 @amp_get_local_size(i32 0) #7
  %67 = lshr i64 %66, 3
  %68 = and i64 %66, 7
  %69 = icmp slt i64 %61, %68
  %70 = zext i1 %69 to i64
  %71 = add nuw nsw i64 %70, %67
  %72 = trunc i64 %71 to i32
  %73 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q1116, i64 0, i32 2, i64 %64
  store i32 %72, i32 addrspace(3)* %73, align 4, !tbaa !26
  br label %74

; <label>:74                                      ; preds = %63, %24
  call spir_func void @hc_barrier(i32 1) #8
  %75 = call spir_func i64 @hc_get_group_id(i32 0) #7
  %76 = shl nsw i64 %75, 9
  %77 = add nsw i64 %76, %61
  %78 = trunc i64 %77 to i32
  %79 = icmp slt i32 %78, %19
  br i1 %79, label %80, label %_Z10visit_nodeiiPN2hc12short_vector5int_2ES2_R11LocalQueuesPiS5_S5_i.exit.i

; <label>:80                                      ; preds = %74
  %sext.i = shl i64 %77, 32
  %81 = ashr exact i64 %sext.i, 32
  %82 = getelementptr inbounds i32 addrspace(1)* %13, i64 %81
  %83 = load i32 addrspace(1)* %82, align 4, !tbaa !26
  %84 = sext i32 %83 to i64
  %85 = getelementptr inbounds i32 addrspace(1)* %17, i64 %84
  store i32 16677221, i32 addrspace(1)* %85, align 4, !tbaa !26
  %86 = getelementptr inbounds i32 addrspace(1)* %18, i64 %84
  %87 = load i32 addrspace(1)* %86, align 4, !tbaa !26
  %88 = getelementptr inbounds %"class.hc::short_vector::int_2.0" addrspace(1)* %15, i64 %84, i32 0
  %89 = load i32 addrspace(1)* %88, align 4, !tbaa !28
  %90 = getelementptr inbounds %"class.hc::short_vector::int_2.0" addrspace(1)* %15, i64 %84, i32 1
  %91 = load i32 addrspace(1)* %90, align 4, !tbaa !30
  %92 = add nsw i32 %91, %89
  %93 = icmp sgt i32 %91, 0
  br i1 %93, label %.lr.ph.i.i, label %_Z10visit_nodeiiPN2hc12short_vector5int_2ES2_R11LocalQueuesPiS5_S5_i.exit.i

.lr.ph.i.i:                                       ; preds = %80
  %94 = and i64 %61, 7
  %95 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q1116, i64 0, i32 0, i64 %94
  %96 = sext i32 %89 to i64
  br label %97

; <label>:97                                      ; preds = %_ZN11LocalQueues6appendEiPii.exit.i.i, %.lr.ph.i.i
  %indvars.iv.i.i = phi i64 [ %96, %.lr.ph.i.i ], [ %indvars.iv.next.i.i, %_ZN11LocalQueues6appendEiPii.exit.i.i ]
  %98 = getelementptr inbounds %"class.hc::short_vector::int_2.0" addrspace(1)* %16, i64 %indvars.iv.i.i, i32 0
  %99 = load i32 addrspace(1)* %98, align 4, !tbaa !28
  %100 = getelementptr inbounds %"class.hc::short_vector::int_2.0" addrspace(1)* %16, i64 %indvars.iv.i.i, i32 1
  %101 = load i32 addrspace(1)* %100, align 4, !tbaa !30
  %102 = add nsw i32 %101, %87
  %103 = sext i32 %99 to i64
  %104 = getelementptr inbounds i32 addrspace(1)* %18, i64 %103
  %105 = call spir_func i32 @atomic_min_int_global(i32 addrspace(1)* %104, i32 %102) #9
  %106 = icmp sgt i32 %105, %102
  br i1 %106, label %107, label %_ZN11LocalQueues6appendEiPii.exit.i.i

; <label>:107                                     ; preds = %97
  %108 = getelementptr inbounds i32 addrspace(1)* %17, i64 %103
  %109 = call spir_func i32 @atomic_exchange_int_global(i32 addrspace(1)* %108, i32 %21) #9
  %110 = icmp eq i32 %109, %21
  br i1 %110, label %_ZN11LocalQueues6appendEiPii.exit.i.i, label %111

; <label>:111                                     ; preds = %107
  %112 = call spir_func i32 @atomic_add_int_local(i32 addrspace(3)* %95, i32 1) #9
  %113 = icmp sgt i32 %112, 399
  br i1 %113, label %114, label %115

; <label>:114                                     ; preds = %111
  store i32 1, i32 addrspace(1)* %23, align 4, !tbaa !26
  br label %_ZN11LocalQueues6appendEiPii.exit.i.i

; <label>:115                                     ; preds = %111
  %116 = sext i32 %112 to i64
  %117 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q1116, i64 0, i32 1, i64 %94, i64 %116
  store i32 %99, i32 addrspace(3)* %117, align 4, !tbaa !26
  br label %_ZN11LocalQueues6appendEiPii.exit.i.i

_ZN11LocalQueues6appendEiPii.exit.i.i:            ; preds = %115, %114, %107, %97
  %indvars.iv.next.i.i = add nsw i64 %indvars.iv.i.i, 1
  %118 = trunc i64 %indvars.iv.next.i.i to i32
  %119 = icmp slt i32 %118, %92
  br i1 %119, label %97, label %_Z10visit_nodeiiPN2hc12short_vector5int_2ES2_R11LocalQueuesPiS5_S5_i.exit.loopexit.i

_Z10visit_nodeiiPN2hc12short_vector5int_2ES2_R11LocalQueuesPiS5_S5_i.exit.loopexit.i: ; preds = %_ZN11LocalQueues6appendEiPii.exit.i.i
  br label %_Z10visit_nodeiiPN2hc12short_vector5int_2ES2_R11LocalQueuesPiS5_S5_i.exit.i

_Z10visit_nodeiiPN2hc12short_vector5int_2ES2_R11LocalQueuesPiS5_S5_i.exit.i: ; preds = %_Z10visit_nodeiiPN2hc12short_vector5int_2ES2_R11LocalQueuesPiS5_S5_i.exit.loopexit.i, %80, %74
  call spir_func void @hc_barrier(i32 1) #8
  %120 = icmp eq i64 %61, 0
  br i1 %120, label %121, label %154

; <label>:121                                     ; preds = %_Z10visit_nodeiiPN2hc12short_vector5int_2ES2_R11LocalQueuesPiS5_S5_i.exit.i
  %122 = getelementptr inbounds [8 x i32] addrspace(3)* @ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E8prefix_q1217, i64 0, i64 0
  store i32 0, i32 addrspace(3)* %122, align 16, !tbaa !26
  %123 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q1116, i64 0, i32 0, i64 0
  %124 = load i32 addrspace(3)* %123, align 4, !tbaa !26
  %125 = getelementptr inbounds [8 x i32] addrspace(3)* @ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E8prefix_q1217, i64 0, i64 1
  store i32 %124, i32 addrspace(3)* %125, align 4, !tbaa !26
  %126 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q1116, i64 0, i32 0, i64 1
  %127 = load i32 addrspace(3)* %126, align 4, !tbaa !26
  %128 = add nsw i32 %127, %124
  %129 = getelementptr inbounds [8 x i32] addrspace(3)* @ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E8prefix_q1217, i64 0, i64 2
  store i32 %128, i32 addrspace(3)* %129, align 8, !tbaa !26
  %130 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q1116, i64 0, i32 0, i64 2
  %131 = load i32 addrspace(3)* %130, align 4, !tbaa !26
  %132 = add nsw i32 %131, %128
  %133 = getelementptr inbounds [8 x i32] addrspace(3)* @ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E8prefix_q1217, i64 0, i64 3
  store i32 %132, i32 addrspace(3)* %133, align 4, !tbaa !26
  %134 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q1116, i64 0, i32 0, i64 3
  %135 = load i32 addrspace(3)* %134, align 4, !tbaa !26
  %136 = add nsw i32 %135, %132
  %137 = getelementptr inbounds [8 x i32] addrspace(3)* @ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E8prefix_q1217, i64 0, i64 4
  store i32 %136, i32 addrspace(3)* %137, align 16, !tbaa !26
  %138 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q1116, i64 0, i32 0, i64 4
  %139 = load i32 addrspace(3)* %138, align 4, !tbaa !26
  %140 = add nsw i32 %139, %136
  %141 = getelementptr inbounds [8 x i32] addrspace(3)* @ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E8prefix_q1217, i64 0, i64 5
  store i32 %140, i32 addrspace(3)* %141, align 4, !tbaa !26
  %142 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q1116, i64 0, i32 0, i64 5
  %143 = load i32 addrspace(3)* %142, align 4, !tbaa !26
  %144 = add nsw i32 %143, %140
  %145 = getelementptr inbounds [8 x i32] addrspace(3)* @ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E8prefix_q1217, i64 0, i64 6
  store i32 %144, i32 addrspace(3)* %145, align 8, !tbaa !26
  %146 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q1116, i64 0, i32 0, i64 6
  %147 = load i32 addrspace(3)* %146, align 4, !tbaa !26
  %148 = add nsw i32 %147, %144
  %149 = getelementptr inbounds [8 x i32] addrspace(3)* @ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E8prefix_q1217, i64 0, i64 7
  store i32 %148, i32 addrspace(3)* %149, align 4, !tbaa !26
  %150 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q1116, i64 0, i32 0, i64 7
  %151 = load i32 addrspace(3)* %150, align 4, !tbaa !26
  %152 = add nsw i32 %151, %148
  %153 = call spir_func i32 @atomic_add_int_global(i32 addrspace(1)* %20, i32 %152) #9
  store i32 %153, i32 addrspace(3)* @ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E5shift1318, align 4, !tbaa !26
  br label %154

; <label>:154                                     ; preds = %121, %_Z10visit_nodeiiPN2hc12short_vector5int_2ES2_R11LocalQueuesPiS5_S5_i.exit.i
  call spir_func void @hc_barrier(i32 1) #8
  %155 = load i32 addrspace(3)* @ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E5shift1318, align 4, !tbaa !26
  %156 = sext i32 %155 to i64
  %157 = and i64 %61, 7
  %158 = lshr i64 %61, 3
  %159 = trunc i64 %158 to i32
  %160 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q1116, i64 0, i32 0, i64 %157
  %161 = load i32 addrspace(3)* %160, align 4, !tbaa !26
  %162 = icmp slt i32 %159, %161
  br i1 %162, label %.lr.ph.i1.i, label %_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_.exit

.lr.ph.i1.i:                                      ; preds = %154
  %163 = getelementptr inbounds [8 x i32] addrspace(3)* @ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E8prefix_q1217, i64 0, i64 %157
  %164 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q1116, i64 0, i32 2, i64 %157
  br label %165

; <label>:165                                     ; preds = %165, %.lr.ph.i1.i
  %local_shift.01.i.i = phi i32 [ %159, %.lr.ph.i1.i ], [ %174, %165 ]
  %166 = sext i32 %local_shift.01.i.i to i64
  %167 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q1116, i64 0, i32 1, i64 %157, i64 %166
  %168 = load i32 addrspace(3)* %167, align 4, !tbaa !26
  %169 = load i32 addrspace(3)* %163, align 4, !tbaa !26
  %170 = add nsw i32 %169, %local_shift.01.i.i
  %171 = sext i32 %170 to i64
  %.sum.i = add nsw i64 %171, %156
  %172 = getelementptr inbounds i32 addrspace(1)* %14, i64 %.sum.i
  store i32 %168, i32 addrspace(1)* %172, align 4, !tbaa !26
  %173 = load i32 addrspace(3)* %164, align 4, !tbaa !26
  %174 = add nsw i32 %173, %local_shift.01.i.i
  %175 = load i32 addrspace(3)* %160, align 4, !tbaa !26
  %176 = icmp slt i32 %174, %175
  br i1 %176, label %165, label %_ZN11LocalQueues11concatenateEPiRA8_i.exit.loopexit.i

_ZN11LocalQueues11concatenateEPiRA8_i.exit.loopexit.i: ; preds = %165
  br label %_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_.exit

_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_.exit: ; preds = %_ZN11LocalQueues11concatenateEPiRA8_i.exit.loopexit.i, %154
  %177 = bitcast %struct.grid_launch_parm* %25 to i8*
  call spir_func void @llvm.lifetime.end(i64 72, i8* %177)
  call spir_func void @llvm.lifetime.end(i64 72, i8* %45)
  ret void
}

declare i32 @atomic_min_int_global(i32 addrspace(1)*, i32) #6

declare i32 @atomic_exchange_int_global(i32 addrspace(1)*, i32) #6

declare i32 @atomic_add_int_local(i32 addrspace(3)*, i32) #6

declare i32 @atomic_add_int_global(i32 addrspace(1)*, i32) #6

; Function Attrs: uwtable
define spir_kernel void @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1(i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32 addrspace(1)*, i32 addrspace(1)*, %"class.hc::short_vector::int_2.0" addrspace(1)*, %"class.hc::short_vector::int_2.0" addrspace(1)*, i32 addrspace(1)*, i32 addrspace(1)*, i32 addrspace(1)*, i32 addrspace(1)*, i32, i32, i32 addrspace(1)*, i32 addrspace(1)*, i32 addrspace(1)*, i32 addrspace(1)*) #5 align 2 {
  %28 = alloca %struct.grid_launch_parm, align 8
  %29 = alloca %struct.grid_launch_parm, align 16
  %.sroa.23 = alloca [20 x i8], align 4
  %30 = insertelement <4 x i32> undef, i32 %0, i32 0
  %31 = insertelement <4 x i32> %30, i32 %1, i32 1
  %32 = insertelement <4 x i32> %31, i32 %2, i32 2
  %33 = insertelement <4 x i32> %32, i32 %3, i32 3
  %34 = insertelement <4 x i32> undef, i32 %4, i32 0
  %35 = insertelement <4 x i32> %34, i32 %5, i32 1
  %36 = tail call spir_func i64 @amp_get_local_id(i32 2) #7
  %37 = tail call spir_func i64 @amp_get_local_id(i32 1) #7
  %38 = tail call spir_func i64 @amp_get_local_id(i32 0) #7
  %39 = trunc i64 %36 to i32
  %40 = trunc i64 %37 to i32
  %41 = trunc i64 %38 to i32
  %42 = tail call spir_func i64 @amp_get_group_id(i32 2) #7
  %43 = tail call spir_func i64 @amp_get_group_id(i32 1) #7
  %44 = tail call spir_func i64 @amp_get_group_id(i32 0) #7
  %45 = trunc i64 %42 to i32
  %46 = trunc i64 %43 to i32
  %47 = trunc i64 %44 to i32
  %48 = bitcast %struct.grid_launch_parm* %29 to i8*
  call spir_func void @llvm.lifetime.start(i64 72, i8* %48)
  %49 = insertelement <4 x i32> %35, i32 %47, i32 2
  %50 = insertelement <4 x i32> %49, i32 %46, i32 3
  %51 = insertelement <4 x i32> undef, i32 %45, i32 0
  %52 = insertelement <4 x i32> %51, i32 %41, i32 1
  %53 = insertelement <4 x i32> %52, i32 %40, i32 2
  %54 = insertelement <4 x i32> %53, i32 %39, i32 3
  %55 = bitcast %struct.grid_launch_parm* %29 to <4 x i32>*
  store <4 x i32> %33, <4 x i32>* %55, align 16
  %56 = getelementptr inbounds %struct.grid_launch_parm* %29, i64 0, i32 1, i32 1
  %57 = bitcast i32* %56 to <4 x i32>*
  store <4 x i32> %50, <4 x i32>* %57, align 8
  %58 = getelementptr inbounds %struct.grid_launch_parm* %29, i64 0, i32 2, i32 2
  %59 = bitcast i32* %58 to <4 x i32>*
  store <4 x i32> %54, <4 x i32>* %59, align 8
  %60 = getelementptr inbounds %struct.grid_launch_parm* %29, i64 0, i32 4
  store i32 %12, i32* %60, align 16
  %61 = getelementptr inbounds i8* %48, i64 52
  %62 = getelementptr inbounds [20 x i8]* %.sroa.23, i64 0, i64 0
  call spir_func void @llvm.memcpy.p0i8.p0i8.i64(i8* %61, i8* %62, i64 20, i32 4, i1 false)
  %63 = bitcast %struct.grid_launch_parm* %28 to i8*
  call spir_func void @llvm.lifetime.start(i64 72, i8* %63)
  %tmp = bitcast %struct.grid_launch_parm* %28 to i8*
  %tmp1 = bitcast %struct.grid_launch_parm* %29 to i8*
  call spir_func void @llvm.memcpy.p0i8.p0i8.i64(i8* %tmp, i8* %tmp1, i64 72, i32 1, i1 false)
  %64 = call spir_func i64 @amp_get_local_id(i32 0) #7
  %65 = icmp eq i64 %64, 0
  br i1 %65, label %66, label %71

; <label>:66                                      ; preds = %27
  store i32 1, i32 addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E8odd_time1035, align 4, !tbaa !26
  %67 = call spir_func i64 @hc_get_group_id(i32 0) #7
  %68 = icmp eq i64 %67, 0
  br i1 %68, label %69, label %71

; <label>:69                                      ; preds = %66
  %70 = load i32 addrspace(1)* %19, align 4, !tbaa !26
  store volatile i32 %70, i32 addrspace(1)* @no_of_nodes_vol27, align 4, !tbaa !26
  br label %71

; <label>:71                                      ; preds = %69, %66, %27
  %72 = call spir_func i32 @atomic_or_int_global(i32 addrspace(1)* %25, i32 0) #9
  %73 = icmp slt i64 %64, 8
  %sext7.i = shl i64 %64, 32
  %74 = ashr exact i64 %sext7.i, 32
  %75 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E7local_q631, i64 0, i32 0, i64 %74
  %76 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E7local_q631, i64 0, i32 2, i64 %74
  %77 = and i64 %64, 7
  %78 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E7local_q631, i64 0, i32 0, i64 %77
  %79 = lshr i64 %64, 3
  %80 = trunc i64 %79 to i32
  %81 = getelementptr inbounds [8 x i32] addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E8prefix_q732, i64 0, i64 %77
  %82 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E7local_q631, i64 0, i32 2, i64 %77
  br label %83

; <label>:83                                      ; preds = %_Z20start_global_barrieri.exit2.i, %71
  %kt.0.i = phi i32 [ %72, %71 ], [ %215, %_Z20start_global_barrieri.exit2.i ]
  %.0.i = phi i32 [ %21, %71 ], [ %.1.i, %_Z20start_global_barrieri.exit2.i ]
  br i1 %73, label %84, label %92

; <label>:84                                      ; preds = %83
  store i32 0, i32 addrspace(3)* %75, align 4, !tbaa !26
  %85 = call spir_func i64 @amp_get_local_size(i32 0) #7
  %86 = lshr i64 %85, 3
  %87 = and i64 %85, 7
  %88 = icmp slt i64 %64, %87
  %89 = zext i1 %88 to i64
  %90 = add nuw nsw i64 %89, %86
  %91 = trunc i64 %90 to i32
  store i32 %91, i32 addrspace(3)* %76, align 4, !tbaa !26
  br label %92

; <label>:92                                      ; preds = %84, %83
  br i1 %65, label %93, label %95

; <label>:93                                      ; preds = %92
  %94 = load volatile i32 addrspace(1)* @no_of_nodes_vol27, align 4, !tbaa !26
  store i32 %94, i32 addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E14no_of_nodes_sm934, align 4, !tbaa !26
  br label %95

; <label>:95                                      ; preds = %93, %92
  call spir_func void @hc_barrier(i32 1) #8
  %96 = call spir_func i64 @hc_get_group_id(i32 0) #7
  %97 = shl nsw i64 %96, 9
  %98 = add nsw i64 %97, %64
  %99 = trunc i64 %98 to i32
  %100 = load i32 addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E14no_of_nodes_sm934, align 4, !tbaa !26
  %101 = icmp slt i32 %99, %100
  br i1 %101, label %102, label %_Z10visit_nodeiiPN2hc12short_vector5int_2ES2_R11LocalQueuesPiS5_S5_i.exit.i

; <label>:102                                     ; preds = %95
  %103 = load i32 addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E8odd_time1035, align 4, !tbaa !26
  %104 = icmp ne i32 %103, 0
  %105 = select i1 %104, i32 addrspace(1)* %13, i32 addrspace(1)* %14
  %sext.i = shl i64 %98, 32
  %106 = ashr exact i64 %sext.i, 32
  %107 = getelementptr inbounds i32 addrspace(1)* %105, i64 %106
  %108 = call spir_func i32 @atomic_or_int_global(i32 addrspace(1)* %107, i32 0) #9
  %109 = sext i32 %108 to i64
  %110 = getelementptr inbounds i32 addrspace(1)* %17, i64 %109
  store i32 16677221, i32 addrspace(1)* %110, align 4, !tbaa !26
  %111 = getelementptr inbounds i32 addrspace(1)* %18, i64 %109
  %112 = load i32 addrspace(1)* %111, align 4, !tbaa !26
  %113 = getelementptr inbounds %"class.hc::short_vector::int_2.0" addrspace(1)* %15, i64 %109, i32 0
  %114 = load i32 addrspace(1)* %113, align 4, !tbaa !28
  %115 = getelementptr inbounds %"class.hc::short_vector::int_2.0" addrspace(1)* %15, i64 %109, i32 1
  %116 = load i32 addrspace(1)* %115, align 4, !tbaa !30
  %117 = add nsw i32 %116, %114
  %118 = icmp sgt i32 %116, 0
  br i1 %118, label %.lr.ph.i.i, label %_Z10visit_nodeiiPN2hc12short_vector5int_2ES2_R11LocalQueuesPiS5_S5_i.exit.i

.lr.ph.i.i:                                       ; preds = %102
  %119 = sext i32 %114 to i64
  br label %120

; <label>:120                                     ; preds = %_ZN11LocalQueues6appendEiPii.exit.i.i, %.lr.ph.i.i
  %indvars.iv.i.i = phi i64 [ %119, %.lr.ph.i.i ], [ %indvars.iv.next.i.i, %_ZN11LocalQueues6appendEiPii.exit.i.i ]
  %121 = getelementptr inbounds %"class.hc::short_vector::int_2.0" addrspace(1)* %16, i64 %indvars.iv.i.i, i32 0
  %122 = load i32 addrspace(1)* %121, align 4, !tbaa !28
  %123 = getelementptr inbounds %"class.hc::short_vector::int_2.0" addrspace(1)* %16, i64 %indvars.iv.i.i, i32 1
  %124 = load i32 addrspace(1)* %123, align 4, !tbaa !30
  %125 = add nsw i32 %124, %112
  %126 = sext i32 %122 to i64
  %127 = getelementptr inbounds i32 addrspace(1)* %18, i64 %126
  %128 = call spir_func i32 @atomic_min_int_global(i32 addrspace(1)* %127, i32 %125) #9
  %129 = icmp sgt i32 %128, %125
  br i1 %129, label %130, label %_ZN11LocalQueues6appendEiPii.exit.i.i

; <label>:130                                     ; preds = %120
  %131 = getelementptr inbounds i32 addrspace(1)* %17, i64 %126
  %132 = call spir_func i32 @atomic_exchange_int_global(i32 addrspace(1)* %131, i32 %.0.i) #9
  %133 = icmp eq i32 %132, %.0.i
  br i1 %133, label %_ZN11LocalQueues6appendEiPii.exit.i.i, label %134

; <label>:134                                     ; preds = %130
  %135 = call spir_func i32 @atomic_add_int_local(i32 addrspace(3)* %78, i32 1) #9
  %136 = icmp sgt i32 %135, 399
  br i1 %136, label %137, label %138

; <label>:137                                     ; preds = %134
  store i32 1, i32 addrspace(1)* %26, align 4, !tbaa !26
  br label %_ZN11LocalQueues6appendEiPii.exit.i.i

; <label>:138                                     ; preds = %134
  %139 = sext i32 %135 to i64
  %140 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E7local_q631, i64 0, i32 1, i64 %77, i64 %139
  store i32 %122, i32 addrspace(3)* %140, align 4, !tbaa !26
  br label %_ZN11LocalQueues6appendEiPii.exit.i.i

_ZN11LocalQueues6appendEiPii.exit.i.i:            ; preds = %138, %137, %130, %120
  %indvars.iv.next.i.i = add nsw i64 %indvars.iv.i.i, 1
  %141 = trunc i64 %indvars.iv.next.i.i to i32
  %142 = icmp slt i32 %141, %117
  br i1 %142, label %120, label %_Z10visit_nodeiiPN2hc12short_vector5int_2ES2_R11LocalQueuesPiS5_S5_i.exit.loopexit.i

_Z10visit_nodeiiPN2hc12short_vector5int_2ES2_R11LocalQueuesPiS5_S5_i.exit.loopexit.i: ; preds = %_ZN11LocalQueues6appendEiPii.exit.i.i
  br label %_Z10visit_nodeiiPN2hc12short_vector5int_2ES2_R11LocalQueuesPiS5_S5_i.exit.i

_Z10visit_nodeiiPN2hc12short_vector5int_2ES2_R11LocalQueuesPiS5_S5_i.exit.i: ; preds = %_Z10visit_nodeiiPN2hc12short_vector5int_2ES2_R11LocalQueuesPiS5_S5_i.exit.loopexit.i, %102, %95
  call spir_func void @hc_barrier(i32 1) #8
  br i1 %65, label %143, label %176

; <label>:143                                     ; preds = %_Z10visit_nodeiiPN2hc12short_vector5int_2ES2_R11LocalQueuesPiS5_S5_i.exit.i
  %144 = getelementptr inbounds [8 x i32] addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E8prefix_q732, i64 0, i64 0
  store i32 0, i32 addrspace(3)* %144, align 16, !tbaa !26
  %145 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E7local_q631, i64 0, i32 0, i64 0
  %146 = load i32 addrspace(3)* %145, align 4, !tbaa !26
  %147 = getelementptr inbounds [8 x i32] addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E8prefix_q732, i64 0, i64 1
  store i32 %146, i32 addrspace(3)* %147, align 4, !tbaa !26
  %148 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E7local_q631, i64 0, i32 0, i64 1
  %149 = load i32 addrspace(3)* %148, align 4, !tbaa !26
  %150 = add nsw i32 %149, %146
  %151 = getelementptr inbounds [8 x i32] addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E8prefix_q732, i64 0, i64 2
  store i32 %150, i32 addrspace(3)* %151, align 8, !tbaa !26
  %152 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E7local_q631, i64 0, i32 0, i64 2
  %153 = load i32 addrspace(3)* %152, align 4, !tbaa !26
  %154 = add nsw i32 %153, %150
  %155 = getelementptr inbounds [8 x i32] addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E8prefix_q732, i64 0, i64 3
  store i32 %154, i32 addrspace(3)* %155, align 4, !tbaa !26
  %156 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E7local_q631, i64 0, i32 0, i64 3
  %157 = load i32 addrspace(3)* %156, align 4, !tbaa !26
  %158 = add nsw i32 %157, %154
  %159 = getelementptr inbounds [8 x i32] addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E8prefix_q732, i64 0, i64 4
  store i32 %158, i32 addrspace(3)* %159, align 16, !tbaa !26
  %160 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E7local_q631, i64 0, i32 0, i64 4
  %161 = load i32 addrspace(3)* %160, align 4, !tbaa !26
  %162 = add nsw i32 %161, %158
  %163 = getelementptr inbounds [8 x i32] addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E8prefix_q732, i64 0, i64 5
  store i32 %162, i32 addrspace(3)* %163, align 4, !tbaa !26
  %164 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E7local_q631, i64 0, i32 0, i64 5
  %165 = load i32 addrspace(3)* %164, align 4, !tbaa !26
  %166 = add nsw i32 %165, %162
  %167 = getelementptr inbounds [8 x i32] addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E8prefix_q732, i64 0, i64 6
  store i32 %166, i32 addrspace(3)* %167, align 8, !tbaa !26
  %168 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E7local_q631, i64 0, i32 0, i64 6
  %169 = load i32 addrspace(3)* %168, align 4, !tbaa !26
  %170 = add nsw i32 %169, %166
  %171 = getelementptr inbounds [8 x i32] addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E8prefix_q732, i64 0, i64 7
  store i32 %170, i32 addrspace(3)* %171, align 4, !tbaa !26
  %172 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E7local_q631, i64 0, i32 0, i64 7
  %173 = load i32 addrspace(3)* %172, align 4, !tbaa !26
  %174 = add nsw i32 %173, %170
  %175 = call spir_func i32 @atomic_add_int_global(i32 addrspace(1)* %20, i32 %174) #9
  store i32 %175, i32 addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E5shift833, align 4, !tbaa !26
  br label %176

; <label>:176                                     ; preds = %143, %_Z10visit_nodeiiPN2hc12short_vector5int_2ES2_R11LocalQueuesPiS5_S5_i.exit.i
  call spir_func void @hc_barrier(i32 1) #8
  %177 = load i32 addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E8odd_time1035, align 4, !tbaa !26
  %178 = icmp ne i32 %177, 0
  %179 = select i1 %178, i32 addrspace(1)* %14, i32 addrspace(1)* %13
  %180 = load i32 addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E5shift833, align 4, !tbaa !26
  %181 = sext i32 %180 to i64
  %182 = load i32 addrspace(3)* %78, align 4, !tbaa !26
  %183 = icmp slt i32 %80, %182
  br i1 %183, label %.lr.ph.i1.preheader.i, label %_ZN11LocalQueues11concatenateEPiRA8_i.exit.i

.lr.ph.i1.preheader.i:                            ; preds = %176
  br label %.lr.ph.i1.i

.lr.ph.i1.i:                                      ; preds = %.lr.ph.i1.i, %.lr.ph.i1.preheader.i
  %local_shift.01.i.i = phi i32 [ %192, %.lr.ph.i1.i ], [ %80, %.lr.ph.i1.preheader.i ]
  %184 = sext i32 %local_shift.01.i.i to i64
  %185 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E7local_q631, i64 0, i32 1, i64 %77, i64 %184
  %186 = load i32 addrspace(3)* %185, align 4, !tbaa !26
  %187 = load i32 addrspace(3)* %81, align 4, !tbaa !26
  %188 = add nsw i32 %187, %local_shift.01.i.i
  %189 = sext i32 %188 to i64
  %.sum.i = add nsw i64 %189, %181
  %190 = getelementptr inbounds i32 addrspace(1)* %179, i64 %.sum.i
  store i32 %186, i32 addrspace(1)* %190, align 4, !tbaa !26
  %191 = load i32 addrspace(3)* %82, align 4, !tbaa !26
  %192 = add nsw i32 %191, %local_shift.01.i.i
  %193 = load i32 addrspace(3)* %78, align 4, !tbaa !26
  %194 = icmp slt i32 %192, %193
  br i1 %194, label %.lr.ph.i1.i, label %_ZN11LocalQueues11concatenateEPiRA8_i.exit.loopexit.i

_ZN11LocalQueues11concatenateEPiRA8_i.exit.loopexit.i: ; preds = %.lr.ph.i1.i
  br label %_ZN11LocalQueues11concatenateEPiRA8_i.exit.i

_ZN11LocalQueues11concatenateEPiRA8_i.exit.i:     ; preds = %_ZN11LocalQueues11concatenateEPiRA8_i.exit.loopexit.i, %176
  br i1 %65, label %195, label %200

; <label>:195                                     ; preds = %_ZN11LocalQueues11concatenateEPiRA8_i.exit.i
  %196 = load i32 addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E8odd_time1035, align 4, !tbaa !26
  %197 = add nsw i32 %196, 1
  %198 = srem i32 %197, 2
  store i32 %198, i32 addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E8odd_time1035, align 4, !tbaa !26
  %199 = icmp eq i32 %.0.i, 16677219
  %..i = select i1 %199, i32 16677220, i32 16677219
  br label %200

; <label>:200                                     ; preds = %195, %_ZN11LocalQueues11concatenateEPiRA8_i.exit.i
  %.1.i = phi i32 [ %.0.i, %_ZN11LocalQueues11concatenateEPiRA8_i.exit.i ], [ %..i, %195 ]
  call spir_func void @hc_barrier(i32 1) #8
  br i1 %65, label %201, label %_Z20start_global_barrieri.exit.i

; <label>:201                                     ; preds = %200
  %202 = call spir_func i32 @atomic_add_int_global(i32 addrspace(1)* @count24, i32 1) #9
  %203 = mul i32 %kt.0.i, 14
  %204 = add i32 %203, 14
  br label %205

; <label>:205                                     ; preds = %205, %201
  %206 = load volatile i32 addrspace(1)* @count24, align 4, !tbaa !26
  %207 = icmp slt i32 %206, %204
  br i1 %207, label %205, label %_Z20start_global_barrieri.exit.loopexit.i

_Z20start_global_barrieri.exit.loopexit.i:        ; preds = %205
  br label %_Z20start_global_barrieri.exit.i

_Z20start_global_barrieri.exit.i:                 ; preds = %_Z20start_global_barrieri.exit.loopexit.i, %200
  call spir_func void @hc_barrier(i32 1) #8
  %208 = or i64 %96, %64
  %brmerge.i = icmp eq i64 %208, 0
  br i1 %brmerge.i, label %209, label %214

; <label>:209                                     ; preds = %_Z20start_global_barrieri.exit.i
  store volatile i32 0, i32 addrspace(1)* @stay_vol29, align 4, !tbaa !26
  %210 = load i32 addrspace(1)* %20, align 4, !tbaa !26
  %.off.i = add i32 %210, -513
  %211 = icmp ult i32 %.off.i, 6655
  br i1 %211, label %212, label %214

; <label>:212                                     ; preds = %209
  store volatile i32 1, i32 addrspace(1)* @stay_vol29, align 4, !tbaa !26
  %213 = load i32 addrspace(1)* %20, align 4, !tbaa !26
  store volatile i32 %213, i32 addrspace(1)* @no_of_nodes_vol27, align 4, !tbaa !26
  store i32 0, i32 addrspace(1)* %20, align 4, !tbaa !26
  br label %214

; <label>:214                                     ; preds = %212, %209, %_Z20start_global_barrieri.exit.i
  %215 = add nsw i32 %kt.0.i, 2
  call spir_func void @hc_barrier(i32 1) #8
  br i1 %65, label %216, label %_Z20start_global_barrieri.exit2.i

; <label>:216                                     ; preds = %214
  %217 = call spir_func i32 @atomic_add_int_global(i32 addrspace(1)* @count24, i32 1) #9
  %218 = mul nsw i32 %215, 14
  br label %219

; <label>:219                                     ; preds = %219, %216
  %220 = load volatile i32 addrspace(1)* @count24, align 4, !tbaa !26
  %221 = icmp slt i32 %220, %218
  br i1 %221, label %219, label %_Z20start_global_barrieri.exit2.loopexit.i

_Z20start_global_barrieri.exit2.loopexit.i:       ; preds = %219
  br label %_Z20start_global_barrieri.exit2.i

_Z20start_global_barrieri.exit2.i:                ; preds = %_Z20start_global_barrieri.exit2.loopexit.i, %214
  call spir_func void @hc_barrier(i32 1) #8
  %222 = load volatile i32 addrspace(1)* @stay_vol29, align 4, !tbaa !26
  %223 = icmp eq i32 %222, 0
  br i1 %223, label %224, label %83

; <label>:224                                     ; preds = %_Z20start_global_barrieri.exit2.i
  %.not4.i = icmp ne i64 %96, 0
  %.not5.i = xor i1 %65, true
  %brmerge6.i = or i1 %.not4.i, %.not5.i
  br i1 %brmerge6.i, label %_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_.exit, label %225

; <label>:225                                     ; preds = %224
  store i32 %215, i32 addrspace(1)* %25, align 4, !tbaa !26
  %226 = load i32 addrspace(3)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1.ZZ26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_E8odd_time1035, align 4, !tbaa !26
  %227 = add nsw i32 %226, 1
  %228 = srem i32 %227, 2
  store i32 %228, i32 addrspace(1)* %23, align 4, !tbaa !26
  %229 = load volatile i32 addrspace(1)* @no_of_nodes_vol27, align 4, !tbaa !26
  store i32 %229, i32 addrspace(1)* %19, align 4, !tbaa !26
  br label %_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_.exit

_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0_.exit: ; preds = %225, %224
  %230 = bitcast %struct.grid_launch_parm* %28 to i8*
  call spir_func void @llvm.lifetime.end(i64 72, i8* %230)
  call spir_func void @llvm.lifetime.end(i64 72, i8* %48)
  ret void
}

declare i32 @atomic_or_int_global(i32 addrspace(1)*, i32) #6

; Function Attrs: uwtable
define spir_kernel void @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_(i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32 addrspace(1)*, i32 addrspace(1)*, %"class.hc::short_vector::int_2.0" addrspace(1)*, %"class.hc::short_vector::int_2.0" addrspace(1)*, i32 addrspace(1)*, i32 addrspace(1)*, i32, i32 addrspace(1)*, i32, i32, i32 addrspace(1)*) #5 align 2 {
  %25 = alloca %struct.grid_launch_parm, align 8
  %26 = alloca %struct.grid_launch_parm, align 16
  %.sroa.23 = alloca [20 x i8], align 4
  %27 = insertelement <4 x i32> undef, i32 %0, i32 0
  %28 = insertelement <4 x i32> %27, i32 %1, i32 1
  %29 = insertelement <4 x i32> %28, i32 %2, i32 2
  %30 = insertelement <4 x i32> %29, i32 %3, i32 3
  %31 = insertelement <4 x i32> undef, i32 %4, i32 0
  %32 = insertelement <4 x i32> %31, i32 %5, i32 1
  %33 = tail call spir_func i64 @amp_get_local_id(i32 2) #7
  %34 = tail call spir_func i64 @amp_get_local_id(i32 1) #7
  %35 = tail call spir_func i64 @amp_get_local_id(i32 0) #7
  %36 = trunc i64 %33 to i32
  %37 = trunc i64 %34 to i32
  %38 = trunc i64 %35 to i32
  %39 = tail call spir_func i64 @amp_get_group_id(i32 2) #7
  %40 = tail call spir_func i64 @amp_get_group_id(i32 1) #7
  %41 = tail call spir_func i64 @amp_get_group_id(i32 0) #7
  %42 = trunc i64 %39 to i32
  %43 = trunc i64 %40 to i32
  %44 = trunc i64 %41 to i32
  %45 = bitcast %struct.grid_launch_parm* %26 to i8*
  call spir_func void @llvm.lifetime.start(i64 72, i8* %45)
  %46 = insertelement <4 x i32> %32, i32 %44, i32 2
  %47 = insertelement <4 x i32> %46, i32 %43, i32 3
  %48 = insertelement <4 x i32> undef, i32 %42, i32 0
  %49 = insertelement <4 x i32> %48, i32 %38, i32 1
  %50 = insertelement <4 x i32> %49, i32 %37, i32 2
  %51 = insertelement <4 x i32> %50, i32 %36, i32 3
  %52 = bitcast %struct.grid_launch_parm* %26 to <4 x i32>*
  store <4 x i32> %30, <4 x i32>* %52, align 16
  %53 = getelementptr inbounds %struct.grid_launch_parm* %26, i64 0, i32 1, i32 1
  %54 = bitcast i32* %53 to <4 x i32>*
  store <4 x i32> %47, <4 x i32>* %54, align 8
  %55 = getelementptr inbounds %struct.grid_launch_parm* %26, i64 0, i32 2, i32 2
  %56 = bitcast i32* %55 to <4 x i32>*
  store <4 x i32> %51, <4 x i32>* %56, align 8
  %57 = getelementptr inbounds %struct.grid_launch_parm* %26, i64 0, i32 4
  store i32 %12, i32* %57, align 16
  %58 = getelementptr inbounds i8* %45, i64 52
  %59 = getelementptr inbounds [20 x i8]* %.sroa.23, i64 0, i64 0
  call spir_func void @llvm.memcpy.p0i8.p0i8.i64(i8* %58, i8* %59, i64 20, i32 4, i1 false)
  %60 = bitcast %struct.grid_launch_parm* %25 to i8*
  call spir_func void @llvm.lifetime.start(i64 72, i8* %60)
  %tmp = bitcast %struct.grid_launch_parm* %25 to i8*
  %tmp1 = bitcast %struct.grid_launch_parm* %26 to i8*
  call spir_func void @llvm.memcpy.p0i8.p0i8.i64(i8* %tmp, i8* %tmp1, i64 72, i32 1, i1 false)
  %61 = call spir_func i64 @amp_get_local_id(i32 0) #7
  %62 = icmp eq i64 %61, 0
  br i1 %62, label %63, label %.preheader.i

; <label>:63                                      ; preds = %24
  store i32 0, i32 addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7tot_sum548, align 4, !tbaa !26
  br label %.preheader.i

.preheader.i:                                     ; preds = %63, %24
  %64 = icmp slt i64 %61, 8
  %sext6.i = shl i64 %61, 32
  %65 = ashr exact i64 %sext6.i, 32
  %66 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q245, i64 0, i32 0, i64 %65
  %67 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q245, i64 0, i32 2, i64 %65
  %68 = and i64 %61, 7
  %69 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q245, i64 0, i32 0, i64 %68
  %70 = lshr i64 %61, 3
  %71 = trunc i64 %70 to i32
  %72 = getelementptr inbounds [8 x i32] addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E8prefix_q346, i64 0, i64 %68
  %73 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q245, i64 0, i32 2, i64 %68
  %sext18.i = shl i64 %70, 32
  %74 = ashr exact i64 %sext18.i, 32
  br label %.outer.i

.outer.i:                                         ; preds = %185, %.preheader.i
  %.01.ph.i = phi i32 [ %21, %.preheader.i ], [ %..i, %185 ]
  %.0.ph.i = phi i32 [ %19, %.preheader.i ], [ %184, %185 ]
  br label %75

; <label>:75                                      ; preds = %_ZN11LocalQueues11concatenateEPiRA8_i.exit.i, %.outer.i
  %.0.i = phi i32 [ %184, %_ZN11LocalQueues11concatenateEPiRA8_i.exit.i ], [ %.0.ph.i, %.outer.i ]
  br i1 %64, label %76, label %84

; <label>:76                                      ; preds = %75
  store i32 0, i32 addrspace(3)* %66, align 4, !tbaa !26
  %77 = call spir_func i64 @amp_get_local_size(i32 0) #7
  %78 = lshr i64 %77, 3
  %79 = and i64 %77, 7
  %80 = icmp slt i64 %61, %79
  %81 = zext i1 %80 to i64
  %82 = add nuw nsw i64 %81, %78
  %83 = trunc i64 %82 to i32
  store i32 %83, i32 addrspace(3)* %67, align 4, !tbaa !26
  br label %84

; <label>:84                                      ; preds = %76, %75
  call spir_func void @hc_barrier(i32 1) #8
  %85 = call spir_func i64 @hc_get_group_id(i32 0) #7
  %86 = shl nsw i64 %85, 9
  %87 = add nsw i64 %86, %61
  %88 = trunc i64 %87 to i32
  %89 = icmp slt i32 %88, %.0.i
  br i1 %89, label %90, label %_Z10visit_nodeiiPN2hc12short_vector5int_2ES2_R11LocalQueuesPiS5_S5_i.exit.i

; <label>:90                                      ; preds = %84
  %91 = load i32 addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7tot_sum548, align 4, !tbaa !26
  %92 = icmp eq i32 %91, 0
  %sext.i = shl i64 %87, 32
  %93 = ashr exact i64 %sext.i, 32
  br i1 %92, label %94, label %96

; <label>:94                                      ; preds = %90
  %95 = getelementptr inbounds i32 addrspace(1)* %13, i64 %93
  br label %98

; <label>:96                                      ; preds = %90
  %97 = getelementptr inbounds [512 x i32] addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7next_wf447, i64 0, i64 %93
  br label %98

; <label>:98                                      ; preds = %96, %94
  %pid.0.in.i = phi i32 addrspace(3)* [ %95, %94 ], [ %97, %96 ]
  %pid.0.i = load i32 addrspace(3)* %pid.0.in.i, align 4
  %99 = sext i32 %pid.0.i to i64
  %100 = getelementptr inbounds i32 addrspace(1)* %17, i64 %99
  store i32 16677221, i32 addrspace(1)* %100, align 4, !tbaa !26
  %101 = getelementptr inbounds i32 addrspace(1)* %18, i64 %99
  %102 = load i32 addrspace(1)* %101, align 4, !tbaa !26
  %103 = getelementptr inbounds %"class.hc::short_vector::int_2.0" addrspace(1)* %15, i64 %99, i32 0
  %104 = load i32 addrspace(1)* %103, align 4, !tbaa !28
  %105 = getelementptr inbounds %"class.hc::short_vector::int_2.0" addrspace(1)* %15, i64 %99, i32 1
  %106 = load i32 addrspace(1)* %105, align 4, !tbaa !30
  %107 = add nsw i32 %106, %104
  %108 = icmp sgt i32 %106, 0
  br i1 %108, label %.lr.ph.i.i, label %_Z10visit_nodeiiPN2hc12short_vector5int_2ES2_R11LocalQueuesPiS5_S5_i.exit.i

.lr.ph.i.i:                                       ; preds = %98
  %109 = sext i32 %104 to i64
  br label %110

; <label>:110                                     ; preds = %_ZN11LocalQueues6appendEiPii.exit.i.i, %.lr.ph.i.i
  %indvars.iv.i.i = phi i64 [ %109, %.lr.ph.i.i ], [ %indvars.iv.next.i.i, %_ZN11LocalQueues6appendEiPii.exit.i.i ]
  %111 = getelementptr inbounds %"class.hc::short_vector::int_2.0" addrspace(1)* %16, i64 %indvars.iv.i.i, i32 0
  %112 = load i32 addrspace(1)* %111, align 4, !tbaa !28
  %113 = getelementptr inbounds %"class.hc::short_vector::int_2.0" addrspace(1)* %16, i64 %indvars.iv.i.i, i32 1
  %114 = load i32 addrspace(1)* %113, align 4, !tbaa !30
  %115 = add nsw i32 %114, %102
  %116 = sext i32 %112 to i64
  %117 = getelementptr inbounds i32 addrspace(1)* %18, i64 %116
  %118 = call spir_func i32 @atomic_min_int_global(i32 addrspace(1)* %117, i32 %115) #9
  %119 = icmp sgt i32 %118, %115
  br i1 %119, label %120, label %_ZN11LocalQueues6appendEiPii.exit.i.i

; <label>:120                                     ; preds = %110
  %121 = getelementptr inbounds i32 addrspace(1)* %17, i64 %116
  %122 = call spir_func i32 @atomic_exchange_int_global(i32 addrspace(1)* %121, i32 %.01.ph.i) #9
  %123 = icmp eq i32 %122, %.01.ph.i
  br i1 %123, label %_ZN11LocalQueues6appendEiPii.exit.i.i, label %124

; <label>:124                                     ; preds = %120
  %125 = call spir_func i32 @atomic_add_int_local(i32 addrspace(3)* %69, i32 1) #9
  %126 = icmp sgt i32 %125, 399
  br i1 %126, label %127, label %128

; <label>:127                                     ; preds = %124
  store i32 1, i32 addrspace(1)* %23, align 4, !tbaa !26
  br label %_ZN11LocalQueues6appendEiPii.exit.i.i

; <label>:128                                     ; preds = %124
  %129 = sext i32 %125 to i64
  %130 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q245, i64 0, i32 1, i64 %68, i64 %129
  store i32 %112, i32 addrspace(3)* %130, align 4, !tbaa !26
  br label %_ZN11LocalQueues6appendEiPii.exit.i.i

_ZN11LocalQueues6appendEiPii.exit.i.i:            ; preds = %128, %127, %120, %110
  %indvars.iv.next.i.i = add nsw i64 %indvars.iv.i.i, 1
  %131 = trunc i64 %indvars.iv.next.i.i to i32
  %132 = icmp slt i32 %131, %107
  br i1 %132, label %110, label %_Z10visit_nodeiiPN2hc12short_vector5int_2ES2_R11LocalQueuesPiS5_S5_i.exit.loopexit.i

_Z10visit_nodeiiPN2hc12short_vector5int_2ES2_R11LocalQueuesPiS5_S5_i.exit.loopexit.i: ; preds = %_ZN11LocalQueues6appendEiPii.exit.i.i
  br label %_Z10visit_nodeiiPN2hc12short_vector5int_2ES2_R11LocalQueuesPiS5_S5_i.exit.i

_Z10visit_nodeiiPN2hc12short_vector5int_2ES2_R11LocalQueuesPiS5_S5_i.exit.i: ; preds = %_Z10visit_nodeiiPN2hc12short_vector5int_2ES2_R11LocalQueuesPiS5_S5_i.exit.loopexit.i, %98, %84
  call spir_func void @hc_barrier(i32 1) #8
  br i1 %62, label %133, label %165

; <label>:133                                     ; preds = %_Z10visit_nodeiiPN2hc12short_vector5int_2ES2_R11LocalQueuesPiS5_S5_i.exit.i
  %134 = getelementptr inbounds [8 x i32] addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E8prefix_q346, i64 0, i64 0
  store i32 0, i32 addrspace(3)* %134, align 16, !tbaa !26
  %135 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q245, i64 0, i32 0, i64 0
  %136 = load i32 addrspace(3)* %135, align 4, !tbaa !26
  %137 = getelementptr inbounds [8 x i32] addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E8prefix_q346, i64 0, i64 1
  store i32 %136, i32 addrspace(3)* %137, align 4, !tbaa !26
  %138 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q245, i64 0, i32 0, i64 1
  %139 = load i32 addrspace(3)* %138, align 4, !tbaa !26
  %140 = add nsw i32 %139, %136
  %141 = getelementptr inbounds [8 x i32] addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E8prefix_q346, i64 0, i64 2
  store i32 %140, i32 addrspace(3)* %141, align 8, !tbaa !26
  %142 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q245, i64 0, i32 0, i64 2
  %143 = load i32 addrspace(3)* %142, align 4, !tbaa !26
  %144 = add nsw i32 %143, %140
  %145 = getelementptr inbounds [8 x i32] addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E8prefix_q346, i64 0, i64 3
  store i32 %144, i32 addrspace(3)* %145, align 4, !tbaa !26
  %146 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q245, i64 0, i32 0, i64 3
  %147 = load i32 addrspace(3)* %146, align 4, !tbaa !26
  %148 = add nsw i32 %147, %144
  %149 = getelementptr inbounds [8 x i32] addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E8prefix_q346, i64 0, i64 4
  store i32 %148, i32 addrspace(3)* %149, align 16, !tbaa !26
  %150 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q245, i64 0, i32 0, i64 4
  %151 = load i32 addrspace(3)* %150, align 4, !tbaa !26
  %152 = add nsw i32 %151, %148
  %153 = getelementptr inbounds [8 x i32] addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E8prefix_q346, i64 0, i64 5
  store i32 %152, i32 addrspace(3)* %153, align 4, !tbaa !26
  %154 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q245, i64 0, i32 0, i64 5
  %155 = load i32 addrspace(3)* %154, align 4, !tbaa !26
  %156 = add nsw i32 %155, %152
  %157 = getelementptr inbounds [8 x i32] addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E8prefix_q346, i64 0, i64 6
  store i32 %156, i32 addrspace(3)* %157, align 8, !tbaa !26
  %158 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q245, i64 0, i32 0, i64 6
  %159 = load i32 addrspace(3)* %158, align 4, !tbaa !26
  %160 = add nsw i32 %159, %156
  %161 = getelementptr inbounds [8 x i32] addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E8prefix_q346, i64 0, i64 7
  store i32 %160, i32 addrspace(3)* %161, align 4, !tbaa !26
  %162 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q245, i64 0, i32 0, i64 7
  %163 = load i32 addrspace(3)* %162, align 4, !tbaa !26
  %164 = add nsw i32 %163, %160
  store i32 %164, i32 addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7tot_sum548, align 4, !tbaa !26
  store i32 %164, i32 addrspace(1)* %20, align 4, !tbaa !26
  br label %165

; <label>:165                                     ; preds = %133, %_Z10visit_nodeiiPN2hc12short_vector5int_2ES2_R11LocalQueuesPiS5_S5_i.exit.i
  call spir_func void @hc_barrier(i32 1) #8
  %166 = load i32 addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7tot_sum548, align 4, !tbaa !26
  %167 = icmp eq i32 %166, 0
  br i1 %167, label %_ZN11LocalQueues11concatenateEPiRA8_i.exit5.loopexit22.i, label %168

; <label>:168                                     ; preds = %165
  %169 = icmp slt i32 %166, 513
  %170 = load i32 addrspace(3)* %69, align 4, !tbaa !26
  %171 = icmp slt i32 %71, %170
  br i1 %169, label %172, label %187

; <label>:172                                     ; preds = %168
  br i1 %171, label %.lr.ph.i2.i, label %_ZN11LocalQueues11concatenateEPiRA8_i.exit.i

.lr.ph.i2.i:                                      ; preds = %172
  %173 = load i32 addrspace(3)* %72, align 4, !tbaa !26
  %174 = load i32 addrspace(3)* %73, align 4, !tbaa !26
  %175 = sext i32 %174 to i64
  %176 = sext i32 %173 to i64
  br label %177

; <label>:177                                     ; preds = %177, %.lr.ph.i2.i
  %indvars.iv.i = phi i64 [ %indvars.iv.next.i, %177 ], [ %74, %.lr.ph.i2.i ]
  %178 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q245, i64 0, i32 1, i64 %68, i64 %indvars.iv.i
  %179 = load i32 addrspace(3)* %178, align 4, !tbaa !26
  %180 = add nsw i64 %176, %indvars.iv.i
  %181 = getelementptr inbounds [512 x i32] addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7next_wf447, i64 0, i64 %180
  store i32 %179, i32 addrspace(3)* %181, align 4, !tbaa !26
  %indvars.iv.next.i = add nsw i64 %indvars.iv.i, %175
  %182 = trunc i64 %indvars.iv.next.i to i32
  %183 = icmp slt i32 %182, %170
  br i1 %183, label %177, label %_ZN11LocalQueues11concatenateEPiRA8_i.exit.loopexit.i

_ZN11LocalQueues11concatenateEPiRA8_i.exit.loopexit.i: ; preds = %177
  br label %_ZN11LocalQueues11concatenateEPiRA8_i.exit.i

_ZN11LocalQueues11concatenateEPiRA8_i.exit.i:     ; preds = %_ZN11LocalQueues11concatenateEPiRA8_i.exit.loopexit.i, %172
  call spir_func void @hc_barrier(i32 1) #8
  %184 = load i32 addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7tot_sum548, align 4, !tbaa !26
  br i1 %62, label %185, label %75

; <label>:185                                     ; preds = %_ZN11LocalQueues11concatenateEPiRA8_i.exit.i
  %186 = icmp eq i32 %.01.ph.i, 16677219
  %..i = select i1 %186, i32 16677220, i32 16677219
  br label %.outer.i

; <label>:187                                     ; preds = %168
  br i1 %171, label %.lr.ph.i3.preheader.i, label %_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_.exit

.lr.ph.i3.preheader.i:                            ; preds = %187
  br label %.lr.ph.i3.i

.lr.ph.i3.i:                                      ; preds = %.lr.ph.i3.i, %.lr.ph.i3.preheader.i
  %local_shift.01.i4.i = phi i32 [ %196, %.lr.ph.i3.i ], [ %71, %.lr.ph.i3.preheader.i ]
  %188 = sext i32 %local_shift.01.i4.i to i64
  %189 = getelementptr inbounds %struct.LocalQueues addrspace(3)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_.ZZ17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_E7local_q245, i64 0, i32 1, i64 %68, i64 %188
  %190 = load i32 addrspace(3)* %189, align 4, !tbaa !26
  %191 = load i32 addrspace(3)* %72, align 4, !tbaa !26
  %192 = add nsw i32 %191, %local_shift.01.i4.i
  %193 = sext i32 %192 to i64
  %194 = getelementptr inbounds i32 addrspace(1)* %14, i64 %193
  store i32 %190, i32 addrspace(1)* %194, align 4, !tbaa !26
  %195 = load i32 addrspace(3)* %73, align 4, !tbaa !26
  %196 = add nsw i32 %195, %local_shift.01.i4.i
  %197 = load i32 addrspace(3)* %69, align 4, !tbaa !26
  %198 = icmp slt i32 %196, %197
  br i1 %198, label %.lr.ph.i3.i, label %_ZN11LocalQueues11concatenateEPiRA8_i.exit5.loopexit.i

_ZN11LocalQueues11concatenateEPiRA8_i.exit5.loopexit.i: ; preds = %.lr.ph.i3.i
  br label %_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_.exit

_ZN11LocalQueues11concatenateEPiRA8_i.exit5.loopexit22.i: ; preds = %165
  br label %_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_.exit

_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0_.exit: ; preds = %_ZN11LocalQueues11concatenateEPiRA8_i.exit5.loopexit22.i, %_ZN11LocalQueues11concatenateEPiRA8_i.exit5.loopexit.i, %187
  %199 = bitcast %struct.grid_launch_parm* %25 to i8*
  call spir_func void @llvm.lifetime.end(i64 72, i8* %199)
  call spir_func void @llvm.lifetime.end(i64 72, i8* %45)
  ret void
}

attributes #0 = { nounwind readnone "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
attributes #1 = { nounwind }
attributes #2 = { noduplicate "less-precise-fpmad"="false" "no-frame-pointer-elim"="true" "no-frame-pointer-elim-non-leaf" "no-infs-fp-math"="true" "no-nans-fp-math"="true" "stack-protector-buffer-size"="8" "unsafe-fp-math"="true" "use-soft-float"="false" }
attributes #3 = { nounwind readnone "less-precise-fpmad"="false" "no-frame-pointer-elim"="true" "no-frame-pointer-elim-non-leaf" "no-infs-fp-math"="true" "no-nans-fp-math"="true" "stack-protector-buffer-size"="8" "unsafe-fp-math"="true" "use-soft-float"="false" }
attributes #4 = { nounwind uwtable "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
attributes #5 = { uwtable "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
attributes #6 = { "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
attributes #7 = { nobuiltin nounwind readnone }
attributes #8 = { nobuiltin noduplicate }
attributes #9 = { nobuiltin }

!llvm.ident = !{!0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !1, !1, !1, !1}
!opencl.kernels = !{!2, !8, !10, !16, !22}

!0 = metadata !{metadata !"HCC clang version 3.5.0  (based on HCC 0.10.16186-d14f969-7461349 LLVM 3.5.0svn)"}
!1 = metadata !{metadata !"HCC clang version 3.5.0  (based on HCC 0.10.16256-042a253-061e735 LLVM 3.5.0svn)"}
!2 = metadata !{void (i64, i8 addrspace(1)*, i8)* @ZZ16ihipMemsetKernelIcEN2hc17completion_futureEP12ihipStream_tPT_S4_mENUlNS0_11tiled_indexILi1EEEE_19__cxxamp_trampolineEmPcc, metadata !3, metadata !4, metadata !5, metadata !6, metadata !7}
!3 = metadata !{metadata !"kernel_arg_addr_space", i32 0, i32 0, i32 0}
!4 = metadata !{metadata !"kernel_arg_access_qual", metadata !"none", metadata !"none", metadata !"none"}
!5 = metadata !{metadata !"kernel_arg_type", metadata !"size_t", metadata !"char*", metadata !"char"}
!6 = metadata !{metadata !"kernel_arg_type_qual", metadata !"", metadata !"", metadata !""}
!7 = metadata !{metadata !"kernel_arg_name", metadata !"", metadata !"", metadata !""}
!8 = metadata !{void (i64, i32 addrspace(1)*, i32)* @ZZ16ihipMemsetKernelIjEN2hc17completion_futureEP12ihipStream_tPT_S4_mENUlNS0_11tiled_indexILi1EEEE_19__cxxamp_trampolineEmPjj, metadata !3, metadata !4, metadata !9, metadata !6, metadata !7}
!9 = metadata !{metadata !"kernel_arg_type", metadata !"size_t", metadata !"uint*", metadata !"uint"}
!10 = metadata !{void (i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32 addrspace(1)*, i32 addrspace(1)*, %"class.hc::short_vector::int_2.0" addrspace(1)*, %"class.hc::short_vector::int_2.0" addrspace(1)*, i32 addrspace(1)*, i32 addrspace(1)*, i32, i32 addrspace(1)*, i32, i32, i32 addrspace(1)*)* @ZN12_GLOBAL__N_189_Z10BFS_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_, metadata !11, metadata !12, metadata !13, metadata !14, metadata !15}
!11 = metadata !{metadata !"kernel_arg_addr_space", i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0}
!12 = metadata !{metadata !"kernel_arg_access_qual", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none"}
!13 = metadata !{metadata !"kernel_arg_type", metadata !"int", metadata !"int", metadata !"int", metadata !"int", metadata !"int", metadata !"int", metadata !"int", metadata !"int", metadata !"int", metadata !"int", metadata !"int", metadata !"int", metadata !"uint", metadata !"int*", metadata !"int*", metadata !"class_hc__short_vector__int_2_gl*", metadata !"class_hc__short_vector__int_2_gl*", metadata !"int*", metadata !"int*", metadata !"int", metadata !"int*", metadata !"int", metadata !"int", metadata !"int*"}
!14 = metadata !{metadata !"kernel_arg_type_qual", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"const", metadata !"", metadata !"const", metadata !"const", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !""}
!15 = metadata !{metadata !"kernel_arg_name", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !""}
!16 = metadata !{void (i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32 addrspace(1)*, i32 addrspace(1)*, %"class.hc::short_vector::int_2.0" addrspace(1)*, %"class.hc::short_vector::int_2.0" addrspace(1)*, i32 addrspace(1)*, i32 addrspace(1)*, i32 addrspace(1)*, i32 addrspace(1)*, i32, i32, i32 addrspace(1)*, i32 addrspace(1)*, i32 addrspace(1)*, i32 addrspace(1)*)* @ZN12_GLOBAL__N_1116_Z26BFS_kernel_multi_blk_inGPU16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_S0_S0_iiS0_S0_S0_S0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPiS1_PK32class_hc__short_vector__int_2_glS4_S1_S1_S1_S1_iiS1_PKiS1_S1, metadata !17, metadata !18, metadata !19, metadata !20, metadata !21}
!17 = metadata !{metadata !"kernel_arg_addr_space", i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0, i32 0}
!18 = metadata !{metadata !"kernel_arg_access_qual", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none", metadata !"none"}
!19 = metadata !{metadata !"kernel_arg_type", metadata !"int", metadata !"int", metadata !"int", metadata !"int", metadata !"int", metadata !"int", metadata !"int", metadata !"int", metadata !"int", metadata !"int", metadata !"int", metadata !"int", metadata !"uint", metadata !"int*", metadata !"int*", metadata !"class_hc__short_vector__int_2_gl*", metadata !"class_hc__short_vector__int_2_gl*", metadata !"int*", metadata !"int*", metadata !"int*", metadata !"int*", metadata !"int", metadata !"int", metadata !"int*", metadata !"int*", metadata !"int*", metadata !"int*"}
!20 = metadata !{metadata !"kernel_arg_type_qual", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"const", metadata !"const", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"const", metadata !"", metadata !""}
!21 = metadata !{metadata !"kernel_arg_name", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !"", metadata !""}
!22 = metadata !{void (i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32, i32 addrspace(1)*, i32 addrspace(1)*, %"class.hc::short_vector::int_2.0" addrspace(1)*, %"class.hc::short_vector::int_2.0" addrspace(1)*, i32 addrspace(1)*, i32 addrspace(1)*, i32, i32 addrspace(1)*, i32, i32, i32 addrspace(1)*)* @ZN12_GLOBAL__N_196_Z17BFS_in_GPU_kernel16grid_launch_parmPiS0_PN2hc12short_vector5int_2ES4_S0_S0_iS0_iiS0__functor19__cxxamp_trampolineEiiiiiiiiiiiijPKiPiPK32class_hc__short_vector__int_2_glS6_S3_S3_iS3_iiS3_, metadata !11, metadata !12, metadata !13, metadata !14, metadata !15}
!23 = metadata !{metadata !24, metadata !24, i64 0}
!24 = metadata !{metadata !"omnipotent char", metadata !25, i64 0}
!25 = metadata !{metadata !"Simple C/C++ TBAA"}
!26 = metadata !{metadata !27, metadata !27, i64 0}
!27 = metadata !{metadata !"int", metadata !24, i64 0}
!28 = metadata !{metadata !29, metadata !27, i64 0}
!29 = metadata !{metadata !"_ZTSN2hc12short_vector5int_2E", metadata !27, i64 0, metadata !27, i64 4}
!30 = metadata !{metadata !29, metadata !27, i64 4}