; LLVM IR from AdaptiveCpp/AdaptiveCpp :: src/runtime/ocl (ocl_queue)

; Unoptimized IR (clang -disable-llvm-passes):
target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-i128:128-f80:128-n8:16:32:64-S128"
target triple = "x86_64-unknown-linux-gnu"

%"class.cl::Context" = type { %"class.cl::detail::Wrapper.51" }
%"class.cl::detail::Wrapper.51" = type { ptr }
%"class.hipsycl::rt::result" = type { %"class.std::unique_ptr" }
%"class.std::unique_ptr" = type { %"struct.std::__uniq_ptr_data" }
%"struct.std::__uniq_ptr_data" = type { %"class.std::__uniq_ptr_impl" }
%"class.std::__uniq_ptr_impl" = type { %"class.std::tuple" }
%"class.std::tuple" = type { %"struct.std::_Tuple_impl" }
%"struct.std::_Tuple_impl" = type { %"struct.std::_Head_base.57" }
%"struct.std::_Head_base.57" = type { ptr }
%"class.cl::UserEvent" = type { %"class.cl::Event" }
%"class.cl::Event" = type { %"class.cl::detail::Wrapper.61" }
%"class.cl::detail::Wrapper.61" = type { ptr }
%"class.hipsycl::rt::source_location" = type <{ %"class.std::__cxx11::basic_string", %"class.std::__cxx11::basic_string", i32, [4 x i8] }>
%"class.std::__cxx11::basic_string" = type { %"struct.std::__cxx11::basic_string<char>::_Alloc_hider", i64, %union.anon }
%"struct.std::__cxx11::basic_string<char>::_Alloc_hider" = type { ptr }
%union.anon = type { i64, [8 x i8] }
%"class.std::allocator.58" = type { i8 }
%"class.hipsycl::rt::error_info" = type <{ %"class.std::__cxx11::basic_string", %"class.hipsycl::rt::error_code", i32, [4 x i8] }>
%"class.hipsycl::rt::error_code" = type { %"class.std::__cxx11::basic_string", i8, i32 }
%"class.std::vector.74" = type { %"struct.std::_Vector_base.75" }
%"struct.std::_Vector_base.75" = type { %"struct.std::_Vector_base<cl::Event, std::allocator<cl::Event>>::_Vector_impl" }
%"struct.std::_Vector_base<cl::Event, std::allocator<cl::Event>>::_Vector_impl" = type { %"struct.std::_Vector_base<cl::Event, std::allocator<cl::Event>>::_Vector_impl_data" }
%"struct.std::_Vector_base<cl::Event, std::allocator<cl::Event>>::_Vector_impl_data" = type { ptr, ptr, ptr }
%"class.std::initializer_list" = type { ptr, i64 }
%"class.std::allocator.76" = type { i8 }
%"class.std::function.116" = type { %"class.std::_Function_base", ptr }
%"class.std::_Function_base" = type { %"union.std::_Any_data", ptr }
%"union.std::_Any_data" = type { %"union.std::_Nocopy_types" }
%"union.std::_Nocopy_types" = type { { i64, i64 } }
%class.anon.118 = type { %"class.cl::UserEvent", %"class.std::shared_ptr.92" }
%"class.std::shared_ptr.92" = type { %"class.std::__shared_ptr.93" }
%"class.std::__shared_ptr.93" = type { ptr, %"class.std::__shared_count" }
%"class.std::__shared_count" = type { ptr }
%"class.hipsycl::rt::ocl_queue" = type { %"class.hipsycl::rt::inorder_queue", ptr, i64, %"class.cl::CommandQueue", %"class.hipsycl::rt::ocl_sscp_code_object_invoker", %"class.hipsycl::rt::worker_thread", %"class.std::shared_ptr", %"struct.hipsycl::rt::ocl_queue::protected_state", %"class.hipsycl::common::spin_lock", %"class.hipsycl::glue::jit::cxx_argument_mapper", %"class.hipsycl::rt::kernel_configuration", %"class.std::unordered_map" }
%"class.hipsycl::rt::inorder_queue" = type { ptr }
%"class.cl::CommandQueue" = type { %"class.cl::detail::Wrapper" }
%"class.cl::detail::Wrapper" = type { ptr }
%"class.hipsycl::rt::ocl_sscp_code_object_invoker" = type { %"class.hipsycl::rt::sscp_code_object_invoker", ptr }
%"class.hipsycl::rt::sscp_code_object_invoker" = type { ptr }
%"class.hipsycl::rt::worker_thread" = type { %"class.std::thread", %"struct.std::atomic", %"class.std::condition_variable", %"class.std::mutex", %"class.std::queue" }
%"class.std::thread" = type { %"class.std::thread::id" }
%"class.std::thread::id" = type { i64 }
%"struct.std::atomic" = type { %"struct.std::__atomic_base" }
%"struct.std::__atomic_base" = type { i8 }
%"class.std::condition_variable" = type { %"class.std::__condvar" }
%"class.std::__condvar" = type { %union.pthread_cond_t }
%union.pthread_cond_t = type { %struct.__pthread_cond_s }
%struct.__pthread_cond_s = type { %union.__atomic_wide_counter, %union.__atomic_wide_counter, [2 x i32], [2 x i32], i32, i32, [2 x i32] }
%union.__atomic_wide_counter = type { i64 }
%"class.std::mutex" = type { %"class.std::__mutex_base" }
%"class.std::__mutex_base" = type { %union.pthread_mutex_t }
%union.pthread_mutex_t = type { %struct.__pthread_mutex_s }
%struct.__pthread_mutex_s = type { i32, i32, i32, i32, i32, i16, i16, %struct.__pthread_internal_list }
%struct.__pthread_internal_list = type { ptr, ptr }
%"class.std::queue" = type { %"class.std::deque" }
%"class.std::deque" = type { %"class.std::_Deque_base" }
%"class.std::_Deque_base" = type { %"struct.std::_Deque_base<std::function<void ()>, std::allocator<std::function<void ()>>>::_Deque_impl" }
%"struct.std::_Deque_base<std::function<void ()>, std::allocator<std::function<void ()>>>::_Deque_impl" = type { %"struct.std::_Deque_base<std::function<void ()>, std::allocator<std::function<void ()>>>::_Deque_impl_data" }
%"struct.std::_Deque_base<std::function<void ()>, std::allocator<std::function<void ()>>>::_Deque_impl_data" = type { ptr, i64, %"struct.std::_Deque_iterator", %"struct.std::_Deque_iterator" }
%"struct.std::_Deque_iterator" = type { ptr, ptr, ptr, ptr }
%"class.std::shared_ptr" = type { %"class.std::__shared_ptr" }
%"class.std::__shared_ptr" = type { ptr, %"class.std::__shared_count" }
%"struct.hipsycl::rt::ocl_queue::protected_state" = type { %"class.std::shared_ptr.0", %"class.std::mutex" }
%"class.std::shared_ptr.0" = type { %"class.std::__shared_ptr.1" }
%"class.std::__shared_ptr.1" = type { ptr, %"class.std::__shared_count" }
%"class.hipsycl::common::spin_lock" = type { %"struct.std::atomic_flag" }
%"struct.std::atomic_flag" = type { %"struct.std::__atomic_flag_base" }
%"struct.std::__atomic_flag_base" = type { i8 }
%"class.hipsycl::glue::jit::cxx_argument_mapper" = type { i8, %"class.std::vector", %"class.std::vector.6" }
%"class.std::vector" = type { %"struct.std::_Vector_base" }
%"struct.std::_Vector_base" = type { %"struct.std::_Vector_base<void *, std::allocator<void *>>::_Vector_impl" }
%"struct.std::_Vector_base<void *, std::allocator<void *>>::_Vector_impl" = type { %"struct.std::_Vector_base<void *, std::allocator<void *>>::_Vector_impl_data" }
%"struct.std::_Vector_base<void *, std::allocator<void *>>::_Vector_impl_data" = type { ptr, ptr, ptr }
%"class.std::vector.6" = type { %"struct.std::_Vector_base.7" }
%"struct.std::_Vector_base.7" = type { %"struct.std::_Vector_base<unsigned long, std::allocator<unsigned long>>::_Vector_impl" }
%"struct.std::_Vector_base<unsigned long, std::allocator<unsigned long>>::_Vector_impl" = type { %"struct.std::_Vector_base<unsigned long, std::allocator<unsigned long>>::_Vector_impl_data" }
%"struct.std::_Vector_base<unsigned long, std::allocator<unsigned long>>::_Vector_impl_data" = type { ptr, ptr, ptr }
%"class.hipsycl::rt::kernel_configuration" = type { %"class.std::vector.11", %"class.std::vector.16", %"class.std::vector.21", %"class.std::vector.26", %"class.std::vector.6", %"class.std::vector.31", %"struct.std::array" }
%"class.std::vector.11" = type { %"struct.std::_Vector_base.12" }
%"struct.std::_Vector_base.12" = type { %"struct.std::_Vector_base<hipsycl::rt::kernel_build_flag, std::allocator<hipsycl::rt::kernel_build_flag>>::_Vector_impl" }
%"struct.std::_Vector_base<hipsycl::rt::kernel_build_flag, std::allocator<hipsycl::rt::kernel_build_flag>>::_Vector_impl" = type { %"struct.std::_Vector_base<hipsycl::rt::kernel_build_flag, std::allocator<hipsycl::rt::kernel_build_flag>>::_Vector_impl_data" }
%"struct.std::_Vector_base<hipsycl::rt::kernel_build_flag, std::allocator<hipsycl::rt::kernel_build_flag>>::_Vector_impl_data" = type { ptr, ptr, ptr }
%"class.std::vector.16" = type { %"struct.std::_Vector_base.17" }
%"struct.std::_Vector_base.17" = type { %"struct.std::_Vector_base<std::pair<hipsycl::rt::kernel_build_option, hipsycl::rt::kernel_configuration::int_or_string>, std::allocator<std::pair<hipsycl::rt::kernel_build_option, hipsycl::rt::kernel_configuration::int_or_string>>>::_Vector_impl" }
%"struct.std::_Vector_base<std::pair<hipsycl::rt::kernel_build_option, hipsycl::rt::kernel_configuration::int_or_string>, std::allocator<std::pair<hipsycl::rt::kernel_build_option, hipsycl::rt::kernel_configuration::int_or_string>>>::_Vector_impl" = type { %"struct.std::_Vector_base<std::pair<hipsycl::rt::kernel_build_option, hipsycl::rt::kernel_configuration::int_or_string>, std::allocator<std::pair<hipsycl::rt::kernel_build_option, hipsycl::rt::kernel_configuration::int_or_string>>>::_Vector_impl_data" }
%"struct.std::_Vector_base<std::pair<hipsycl::rt::kernel_build_option, hipsycl::rt::kernel_configuration::int_or_string>, std::allocator<std::pair<hipsycl::rt::kernel_build_option, hipsycl::rt::kernel_configuration::int_or_string>>>::_Vector_impl_data" = type { ptr, ptr, ptr }
%"class.std::vector.21" = type { %"struct.std::_Vector_base.22" }
%"struct.std::_Vector_base.22" = type { %"struct.std::_Vector_base<std::pair<int, unsigned long>, std::allocator<std::pair<int, unsigned long>>>::_Vector_impl" }
%"struct.std::_Vector_base<std::pair<int, unsigned long>, std::allocator<std::pair<int, unsigned long>>>::_Vector_impl" = type { %"struct.std::_Vector_base<std::pair<int, unsigned long>, std::allocator<std::pair<int, unsigned long>>>::_Vector_impl_data" }
%"struct.std::_Vector_base<std::pair<int, unsigned long>, std::allocator<std::pair<int, unsigned long>>>::_Vector_impl_data" = type { ptr, ptr, ptr }
%"class.std::vector.26" = type { %"struct.std::_Vector_base.27" }
%"struct.std::_Vector_base.27" = type { %"struct.std::_Vector_base<hipsycl::glue::sscp::__acpp_sscp_emit_param_type_annotation_fcall_specialized_config<const hipsycl::glue::sscp::fcall_specialized_config *>, std::allocator<hipsycl::glue::sscp::__acpp_sscp_emit_param_type_annotation_fcall_specialized_config<const hipsycl::glue::sscp::fcall_specialized_config *>>>::_Vector_impl" }
%"struct.std::_Vector_base<hipsycl::glue::sscp::__acpp_sscp_emit_param_type_annotation_fcall_specialized_config<const hipsycl::glue::sscp::fcall_specialized_config *>, std::allocator<hipsycl::glue::sscp::__acpp_sscp_emit_param_type_annotation_fcall_specialized_config<const hipsycl::glue::sscp::fcall_specialized_config *>>>::_Vector_impl" = type { %"struct.std::_Vector_base<hipsycl::glue::sscp::__acpp_sscp_emit_param_type_annotation_fcall_specialized_config<const hipsycl::glue::sscp::fcall_specialized_config *>, std::allocator<hipsycl::glue::sscp::__acpp_sscp_emit_param_type_annotation_fcall_specialized_config<const hipsycl::glue::sscp::fcall_specialized_config *>>>::_Vector_impl_data" }
%"struct.std::_Vector_base<hipsycl::glue::sscp::__acpp_sscp_emit_param_type_annotation_fcall_specialized_config<const hipsycl::glue::sscp::fcall_specialized_config *>, std::allocator<hipsycl::glue::sscp::__acpp_sscp_emit_param_type_annotation_fcall_specialized_config<const hipsycl::glue::sscp::fcall_specialized_config *>>>::_Vector_impl_data" = type { ptr, ptr, ptr }
%"class.std::vector.31" = type { %"struct.std::_Vector_base.32" }
%"struct.std::_Vector_base.32" = type { %"struct.std::_Vector_base<std::pair<int, int>, std::allocator<std::pair<int, int>>>::_Vector_impl" }
%"struct.std::_Vector_base<std::pair<int, int>, std::allocator<std::pair<int, int>>>::_Vector_impl" = type { %"struct.std::_Vector_base<std::pair<int, int>, std::allocator<std::pair<int, int>>>::_Vector_impl_data" }
%"struct.std::_Vector_base<std::pair<int, int>, std::allocator<std::pair<int, int>>>::_Vector_impl_data" = type { ptr, ptr, ptr }
%"struct.std::array" = type { [2 x i64] }
%"class.std::unordered_map" = type { %"class.std::_Hashtable" }
%"class.std::_Hashtable" = type { ptr, i64, %"struct.std::__detail::_Hash_node_base", i64, %"struct.std::__detail::_Prime_rehash_policy", ptr }
%"struct.std::__detail::_Hash_node_base" = type { ptr }
%"struct.std::__detail::_Prime_rehash_policy" = type { float, i64 }

@.str = external hidden unnamed_addr constant [87 x i8], align 1
@.str.3 = external hidden unnamed_addr constant [3 x i8], align 1
@__func__._ZN7hipsycl2rt9ocl_queue24submit_external_wait_forERKSt10shared_ptrINS0_8dag_nodeEE = external hidden unnamed_addr constant [25 x i8], align 1
@.str.17 = external hidden unnamed_addr constant [45 x i8], align 1

declare i32 @__gxx_personality_v0(...)

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

declare void @_ZNK7hipsycl2rt20ocl_hardware_context14get_cl_contextEv(ptr dead_on_unwind writable sret(%"class.cl::Context") align 8, ptr noundef nonnull align 8 dereferenceable(89)) #1

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

; Function Attrs: mustprogress nounwind uwtable
declare void @_ZNSaIcEC2Ev(ptr noundef nonnull align 1 dereferenceable(1)) unnamed_addr #2 align 2

; Function Attrs: mustprogress uwtable
declare dso_local void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_(ptr noundef nonnull align 8 dereferenceable(32), ptr noundef, ptr noundef nonnull align 1 dereferenceable(1)) unnamed_addr #3 align 2

; Function Attrs: mustprogress uwtable
declare dso_local void @_ZN7hipsycl2rt15source_locationC2ERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEES9_i(ptr noundef nonnull align 8 dereferenceable(68), ptr noundef nonnull align 8 dereferenceable(32), ptr noundef nonnull align 8 dereferenceable(32), i32 noundef) unnamed_addr #3 align 2

; Function Attrs: mustprogress uwtable
declare dso_local void @_ZN7hipsycl2rt10error_codeC2ERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEi(ptr noundef nonnull align 8 dereferenceable(40), ptr noundef nonnull align 8 dereferenceable(32), i32 noundef) unnamed_addr #3 align 2

; Function Attrs: mustprogress uwtable
declare dso_local void @_ZN7hipsycl2rt10error_infoC2ERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEENS0_10error_codeENS0_10error_typeE(ptr noundef nonnull align 8 dereferenceable(76), ptr noundef nonnull align 8 dereferenceable(32), ptr noundef, i32 noundef) unnamed_addr #3 align 2

; Function Attrs: inlinehint mustprogress nounwind uwtable
declare dso_local void @_ZN7hipsycl2rt10error_infoD2Ev(ptr noundef nonnull align 8 dereferenceable(76)) unnamed_addr #4 align 2

; Function Attrs: inlinehint mustprogress nounwind uwtable
declare dso_local void @_ZN7hipsycl2rt10error_codeD2Ev(ptr noundef nonnull align 8 dereferenceable(40)) unnamed_addr #4 align 2

; Function Attrs: mustprogress nounwind uwtable
declare void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32)) unnamed_addr #2 align 2

; Function Attrs: mustprogress nounwind uwtable
declare dso_local void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1)) unnamed_addr #2 align 2

; Function Attrs: inlinehint mustprogress nounwind uwtable
declare dso_local void @_ZN7hipsycl2rt15source_locationD2Ev(ptr noundef nonnull align 8 dereferenceable(68)) unnamed_addr #4 align 2

; Function Attrs: mustprogress nounwind uwtable
declare dso_local void @_ZN2cl6detail7WrapperIP11_cl_contextED2Ev(ptr noundef nonnull align 8 dereferenceable(8)) unnamed_addr #2 align 2

; Function Attrs: mustprogress uwtable
declare dso_local void @_ZN2cl5EventC2Ev(ptr noundef nonnull align 8 dereferenceable(8)) unnamed_addr #3 align 2

; Function Attrs: mustprogress uwtable
declare dso_local noundef i32 @_ZNK2cl12CommandQueue26enqueueBarrierWithWaitListEPKSt6vectorINS_5EventESaIS2_EEPS2_(ptr noundef nonnull align 8 dereferenceable(8), ptr noundef, ptr noundef) #3 align 2

; Function Attrs: mustprogress uwtable
declare dso_local void @_ZN7hipsycl2rt9ocl_queue21register_submitted_opEN2cl5EventE(ptr noundef nonnull align 8 dereferenceable(584), ptr noundef) #3 align 2

; Function Attrs: inlinehint mustprogress uwtable
declare dso_local void @_ZN2cl5EventC2ERKS0_(ptr noundef nonnull align 8 dereferenceable(8), ptr noundef nonnull align 8 dereferenceable(8)) unnamed_addr #5 align 2

; Function Attrs: mustprogress nounwind uwtable
declare dso_local void @_ZNSt6vectorIN2cl5EventESaIS1_EED2Ev(ptr noundef nonnull align 8 dereferenceable(24)) unnamed_addr #2 align 2

; Function Attrs: inlinehint mustprogress uwtable
declare dso_local void @_ZN7hipsycl2rt10make_errorERKNS0_15source_locationERKNS0_10error_infoE(ptr dead_on_unwind noalias writable sret(%"class.hipsycl::rt::result") align 8, ptr noundef nonnull align 8 dereferenceable(68), ptr noundef nonnull align 8 dereferenceable(76)) #5

; Function Attrs: inlinehint mustprogress nounwind uwtable
declare dso_local void @_ZN7hipsycl2rt12make_successEv(ptr dead_on_unwind noalias writable sret(%"class.hipsycl::rt::result") align 8) #4

; Function Attrs: mustprogress nounwind uwtable
declare dso_local void @_ZNSaIN2cl5EventEEC2Ev(ptr noundef nonnull align 1 dereferenceable(1)) unnamed_addr #2 align 2

; Function Attrs: mustprogress uwtable
declare dso_local void @_ZNSt6vectorIN2cl5EventESaIS1_EEC2ESt16initializer_listIS1_ERKS2_(ptr noundef nonnull align 8 dereferenceable(24), ptr, i64, ptr noundef nonnull align 1 dereferenceable(1)) unnamed_addr #3 align 2

; Function Attrs: mustprogress uwtable
define dso_local void @_ZN7hipsycl2rt9ocl_queue24submit_external_wait_forERKSt10shared_ptrINS0_8dag_nodeEE(ptr dead_on_unwind noalias writable sret(%"class.hipsycl::rt::result") align 8 %0, ptr noundef nonnull align 8 dereferenceable(584) %1, ptr noundef nonnull align 8 dereferenceable(16) %2) unnamed_addr #3 align 2 personality ptr @__gxx_personality_v0 {
  %4 = alloca ptr, align 8
  %5 = alloca ptr, align 8
  %6 = alloca ptr, align 8
  %7 = alloca ptr, align 8
  %8 = alloca i32, align 4
  %9 = alloca %"class.cl::UserEvent", align 8
  %10 = alloca %"class.cl::Context", align 8
  %11 = alloca ptr, align 8
  %12 = alloca i32, align 4
  %13 = alloca %"class.hipsycl::rt::source_location", align 8
  %14 = alloca %"class.std::__cxx11::basic_string", align 8
  %15 = alloca %"class.std::allocator.58", align 1
  %16 = alloca %"class.std::__cxx11::basic_string", align 8
  %17 = alloca %"class.std::allocator.58", align 1
  %18 = alloca %"class.hipsycl::rt::error_info", align 8
  %19 = alloca %"class.std::__cxx11::basic_string", align 8
  %20 = alloca %"class.std::allocator.58", align 1
  %21 = alloca %"class.hipsycl::rt::error_code", align 8
  %22 = alloca %"class.std::__cxx11::basic_string", align 8
  %23 = alloca %"class.std::allocator.58", align 1
  %24 = alloca i32, align 4
  %25 = alloca %"class.cl::Event", align 8
  %26 = alloca %"class.std::vector.74", align 8
  %27 = alloca %"class.std::initializer_list", align 8
  %28 = alloca [1 x %"class.cl::Event"], align 8
  %29 = alloca ptr, align 8
  %30 = alloca %"class.std::allocator.76", align 1
  %31 = alloca %"class.cl::Event", align 8
  %32 = alloca %"class.std::function.116", align 8
  %33 = alloca %class.anon.118, align 8
  store ptr %0, ptr %4, align 8
  store ptr %1, ptr %5, align 8, !tbaa !5
  store ptr %2, ptr %6, align 8, !tbaa !5
  %34 = load ptr, ptr %5, align 8
  call void @llvm.lifetime.start.p0(i64 8, ptr %7) #6
  %35 = getelementptr inbounds %"class.hipsycl::rt::ocl_queue", ptr %34, i32 0, i32 1
  %36 = load ptr, ptr %35, align 8, !tbaa !9
  %37 = getelementptr inbounds %"class.hipsycl::rt::ocl_queue", ptr %34, i32 0, i32 2
  %38 = load i64, ptr %37, align 8, !tbaa !78
  %39 = load ptr, ptr %36, align 8, !tbaa !79
  %40 = getelementptr inbounds ptr, ptr %39, i64 2
  %41 = load ptr, ptr %40, align 8
  %42 = call noundef ptr %41(ptr noundef nonnull align 8 dereferenceable(60) %36, i64 noundef %38)
  store ptr %42, ptr %7, align 8, !tbaa !5
  call void @llvm.lifetime.start.p0(i64 4, ptr %8) #6
  call void @llvm.lifetime.start.p0(i64 8, ptr %9) #6
  call void @llvm.lifetime.start.p0(i64 8, ptr %10) #6
  %43 = load ptr, ptr %7, align 8, !tbaa !5
  call void @_ZNK7hipsycl2rt20ocl_hardware_context14get_cl_contextEv(ptr dead_on_unwind writable sret(%"class.cl::Context") align 8 %10, ptr noundef nonnull align 8 dereferenceable(89) %43)
  invoke void @_ZN2cl9UserEventC2ERKNS_7ContextEPi(ptr noundef nonnull align 8 dereferenceable(8) %9, ptr noundef nonnull align 8 dereferenceable(8) %10, ptr noundef %8)
          to label %44 unwind label %57

44:                                               ; preds = %3
  call void @_ZN2cl6detail7WrapperIP11_cl_contextED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %10) #6
  call void @llvm.lifetime.end.p0(i64 8, ptr %10) #6
  %45 = load i32, ptr %8, align 4, !tbaa !81
  %46 = icmp ne i32 %45, 0
  br i1 %46, label %47, label %100

47:                                               ; preds = %44
  call void @llvm.lifetime.start.p0(i64 72, ptr %13) #6
  call void @llvm.lifetime.start.p0(i64 32, ptr %14) #6
  call void @llvm.lifetime.start.p0(i64 1, ptr %15) #6
  call void @_ZNSaIcEC2Ev(ptr noundef nonnull align 1 dereferenceable(1) %15) #6
  invoke void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_(ptr noundef nonnull align 8 dereferenceable(32) %14, ptr noundef @__func__._ZN7hipsycl2rt9ocl_queue24submit_external_wait_forERKSt10shared_ptrINS0_8dag_nodeEE, ptr noundef nonnull align 1 dereferenceable(1) %15)
          to label %48 unwind label %61

48:                                               ; preds = %47
  call void @llvm.lifetime.start.p0(i64 32, ptr %16) #6
  call void @llvm.lifetime.start.p0(i64 1, ptr %17) #6
  call void @_ZNSaIcEC2Ev(ptr noundef nonnull align 1 dereferenceable(1) %17) #6
  invoke void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_(ptr noundef nonnull align 8 dereferenceable(32) %16, ptr noundef @.str, ptr noundef nonnull align 1 dereferenceable(1) %17)
          to label %49 unwind label %65

49:                                               ; preds = %48
  invoke void @_ZN7hipsycl2rt15source_locationC2ERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEES9_i(ptr noundef nonnull align 8 dereferenceable(68) %13, ptr noundef nonnull align 8 dereferenceable(32) %14, ptr noundef nonnull align 8 dereferenceable(32) %16, i32 noundef 375)
          to label %50 unwind label %69

50:                                               ; preds = %49
  call void @llvm.lifetime.start.p0(i64 80, ptr %18) #6
  call void @llvm.lifetime.start.p0(i64 32, ptr %19) #6
  call void @llvm.lifetime.start.p0(i64 1, ptr %20) #6
  call void @_ZNSaIcEC2Ev(ptr noundef nonnull align 1 dereferenceable(1) %20) #6
  invoke void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_(ptr noundef nonnull align 8 dereferenceable(32) %19, ptr noundef @.str.17, ptr noundef nonnull align 1 dereferenceable(1) %20)
          to label %51 unwind label %73

51:                                               ; preds = %50
  call void @llvm.lifetime.start.p0(i64 32, ptr %22) #6
  call void @llvm.lifetime.start.p0(i64 1, ptr %23) #6
  call void @_ZNSaIcEC2Ev(ptr noundef nonnull align 1 dereferenceable(1) %23) #6
  invoke void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_(ptr noundef nonnull align 8 dereferenceable(32) %22, ptr noundef @.str.3, ptr noundef nonnull align 1 dereferenceable(1) %23)
          to label %52 unwind label %77

52:                                               ; preds = %51
  %53 = load i32, ptr %8, align 4, !tbaa !81
  invoke void @_ZN7hipsycl2rt10error_codeC2ERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEi(ptr noundef nonnull align 8 dereferenceable(40) %21, ptr noundef nonnull align 8 dereferenceable(32) %22, i32 noundef %53)
          to label %54 unwind label %81

54:                                               ; preds = %52
  invoke void @_ZN7hipsycl2rt10error_infoC2ERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEENS0_10error_codeENS0_10error_typeE(ptr noundef nonnull align 8 dereferenceable(76) %18, ptr noundef nonnull align 8 dereferenceable(32) %19, ptr noundef %21, i32 noundef 1)
          to label %55 unwind label %85

55:                                               ; preds = %54
  invoke void @_ZN7hipsycl2rt10make_errorERKNS0_15source_locationERKNS0_10error_infoE(ptr dead_on_unwind writable sret(%"class.hipsycl::rt::result") align 8 %0, ptr noundef nonnull align 8 dereferenceable(68) %13, ptr noundef nonnull align 8 dereferenceable(76) %18)
          to label %56 unwind label %89

56:                                               ; preds = %55
  call void @_ZN7hipsycl2rt10error_infoD2Ev(ptr noundef nonnull align 8 dereferenceable(76) %18) #6
  call void @_ZN7hipsycl2rt10error_codeD2Ev(ptr noundef nonnull align 8 dereferenceable(40) %21) #6
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %22) #6
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %23) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr %23) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr %22) #6
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %19) #6
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %20) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr %20) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr %19) #6
  call void @llvm.lifetime.end.p0(i64 80, ptr %18) #6
  call void @_ZN7hipsycl2rt15source_locationD2Ev(ptr noundef nonnull align 8 dereferenceable(68) %13) #6
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %16) #6
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %17) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr %17) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr %16) #6
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %14) #6
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %15) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr %15) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr %14) #6
  call void @llvm.lifetime.end.p0(i64 72, ptr %13) #6
  store i32 1, ptr %24, align 4
  br label %182

57:                                               ; preds = %3
  %58 = landingpad { ptr, i32 }
          cleanup
  %59 = extractvalue { ptr, i32 } %58, 0
  store ptr %59, ptr %11, align 8
  %60 = extractvalue { ptr, i32 } %58, 1
  store i32 %60, ptr %12, align 4
  call void @_ZN2cl6detail7WrapperIP11_cl_contextED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %10) #6
  call void @llvm.lifetime.end.p0(i64 8, ptr %10) #6
  br label %184

61:                                               ; preds = %47
  %62 = landingpad { ptr, i32 }
          cleanup
  %63 = extractvalue { ptr, i32 } %62, 0
  store ptr %63, ptr %11, align 8
  %64 = extractvalue { ptr, i32 } %62, 1
  store i32 %64, ptr %12, align 4
  br label %99

65:                                               ; preds = %48
  %66 = landingpad { ptr, i32 }
          cleanup
  %67 = extractvalue { ptr, i32 } %66, 0
  store ptr %67, ptr %11, align 8
  %68 = extractvalue { ptr, i32 } %66, 1
  store i32 %68, ptr %12, align 4
  br label %98

69:                                               ; preds = %49
  %70 = landingpad { ptr, i32 }
          cleanup
  %71 = extractvalue { ptr, i32 } %70, 0
  store ptr %71, ptr %11, align 8
  %72 = extractvalue { ptr, i32 } %70, 1
  store i32 %72, ptr %12, align 4
  br label %97

73:                                               ; preds = %50
  %74 = landingpad { ptr, i32 }
          cleanup
  %75 = extractvalue { ptr, i32 } %74, 0
  store ptr %75, ptr %11, align 8
  %76 = extractvalue { ptr, i32 } %74, 1
  store i32 %76, ptr %12, align 4
  br label %96

77:                                               ; preds = %51
  %78 = landingpad { ptr, i32 }
          cleanup
  %79 = extractvalue { ptr, i32 } %78, 0
  store ptr %79, ptr %11, align 8
  %80 = extractvalue { ptr, i32 } %78, 1
  store i32 %80, ptr %12, align 4
  br label %95

81:                                               ; preds = %52
  %82 = landingpad { ptr, i32 }
          cleanup
  %83 = extractvalue { ptr, i32 } %82, 0
  store ptr %83, ptr %11, align 8
  %84 = extractvalue { ptr, i32 } %82, 1
  store i32 %84, ptr %12, align 4
  br label %94

85:                                               ; preds = %54
  %86 = landingpad { ptr, i32 }
          cleanup
  %87 = extractvalue { ptr, i32 } %86, 0
  store ptr %87, ptr %11, align 8
  %88 = extractvalue { ptr, i32 } %86, 1
  store i32 %88, ptr %12, align 4
  br label %93

89:                                               ; preds = %55
  %90 = landingpad { ptr, i32 }
          cleanup
  %91 = extractvalue { ptr, i32 } %90, 0
  store ptr %91, ptr %11, align 8
  %92 = extractvalue { ptr, i32 } %90, 1
  store i32 %92, ptr %12, align 4
  call void @_ZN7hipsycl2rt10error_infoD2Ev(ptr noundef nonnull align 8 dereferenceable(76) %18) #6
  br label %93

93:                                               ; preds = %89, %85
  call void @_ZN7hipsycl2rt10error_codeD2Ev(ptr noundef nonnull align 8 dereferenceable(40) %21) #6
  br label %94

94:                                               ; preds = %93, %81
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %22) #6
  br label %95

95:                                               ; preds = %94, %77
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %23) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr %23) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr %22) #6
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %19) #6
  br label %96

96:                                               ; preds = %95, %73
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %20) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr %20) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr %19) #6
  call void @llvm.lifetime.end.p0(i64 80, ptr %18) #6
  call void @_ZN7hipsycl2rt15source_locationD2Ev(ptr noundef nonnull align 8 dereferenceable(68) %13) #6
  br label %97

97:                                               ; preds = %96, %69
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %16) #6
  br label %98

98:                                               ; preds = %97, %65
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %17) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr %17) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr %16) #6
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %14) #6
  br label %99

99:                                               ; preds = %98, %61
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %15) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr %15) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr %14) #6
  call void @llvm.lifetime.end.p0(i64 72, ptr %13) #6
  br label %183

100:                                              ; preds = %44
  call void @llvm.lifetime.start.p0(i64 8, ptr %25) #6
  invoke void @_ZN2cl5EventC2Ev(ptr noundef nonnull align 8 dereferenceable(8) %25)
          to label %101 unwind label %130

101:                                              ; preds = %100
  call void @llvm.lifetime.start.p0(i64 24, ptr %26) #6
  call void @llvm.lifetime.start.p0(i64 8, ptr %28) #6
  store ptr %28, ptr %29, align 8
  invoke void @_ZN2cl5EventC2ERKS0_(ptr noundef nonnull align 8 dereferenceable(8) %28, ptr noundef nonnull align 8 dereferenceable(8) %9)
          to label %102 unwind label %134

102:                                              ; preds = %101
  %103 = getelementptr inbounds %"class.std::initializer_list", ptr %27, i32 0, i32 0
  store ptr %28, ptr %103, align 8, !tbaa !83
  %104 = getelementptr inbounds %"class.std::initializer_list", ptr %27, i32 0, i32 1
  store i64 1, ptr %104, align 8, !tbaa !85
  call void @llvm.lifetime.start.p0(i64 1, ptr %30) #6
  call void @_ZNSaIN2cl5EventEEC2Ev(ptr noundef nonnull align 1 dereferenceable(1) %30) #6
  %105 = getelementptr inbounds { ptr, i64 }, ptr %27, i32 0, i32 0
  %106 = load ptr, ptr %105, align 8
  %107 = getelementptr inbounds { ptr, i64 }, ptr %27, i32 0, i32 1
  %108 = load i64, ptr %107, align 8
  invoke void @_ZNSt6vectorIN2cl5EventESaIS1_EEC2ESt16initializer_listIS1_ERKS2_(ptr noundef nonnull align 8 dereferenceable(24) %26, ptr %106, i64 %108, ptr noundef nonnull align 1 dereferenceable(1) %30)
          to label %109 unwind label %145

109:                                              ; preds = %102
  call void @_ZNSt15__new_allocatorIN2cl5EventEED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %30) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr %30) #6
  %110 = getelementptr inbounds [1 x %"class.cl::Event"], ptr %28, i32 0, i32 0
  %111 = getelementptr inbounds %"class.cl::Event", ptr %110, i64 1
  br label %112

112:                                              ; preds = %112, %109
  %113 = phi ptr [ %111, %109 ], [ %114, %112 ]
  %114 = getelementptr inbounds %"class.cl::Event", ptr %113, i64 -1
  call void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %114) #6
  %115 = icmp eq ptr %114, %110
  br i1 %115, label %116, label %112

116:                                              ; preds = %112
  call void @llvm.lifetime.end.p0(i64 8, ptr %28) #6
  %117 = getelementptr inbounds %"class.hipsycl::rt::ocl_queue", ptr %34, i32 0, i32 3
  %118 = invoke noundef i32 @_ZNK2cl12CommandQueue26enqueueBarrierWithWaitListEPKSt6vectorINS_5EventESaIS2_EEPS2_(ptr noundef nonnull align 8 dereferenceable(8) %117, ptr noundef %26, ptr noundef %25)
          to label %119 unwind label %157

119:                                              ; preds = %116
  invoke void @_ZN2cl5EventC2ERKS0_(ptr noundef nonnull align 8 dereferenceable(8) %31, ptr noundef nonnull align 8 dereferenceable(8) %25)
          to label %120 unwind label %157

120:                                              ; preds = %119
  invoke void @_ZN7hipsycl2rt9ocl_queue21register_submitted_opEN2cl5EventE(ptr noundef nonnull align 8 dereferenceable(584) %34, ptr noundef %31)
          to label %121 unwind label %161

121:                                              ; preds = %120
  call void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %31) #6
  %122 = getelementptr inbounds %"class.hipsycl::rt::ocl_queue", ptr %34, i32 0, i32 5
  call void @llvm.lifetime.start.p0(i64 24, ptr %33) #6
  %123 = getelementptr inbounds %class.anon.118, ptr %33, i32 0, i32 0
  invoke void @_ZN2cl9UserEventC2ERKS0_(ptr noundef nonnull align 8 dereferenceable(8) %123, ptr noundef nonnull align 8 dereferenceable(8) %9)
          to label %124 unwind label %165

124:                                              ; preds = %121
  %125 = getelementptr inbounds %class.anon.118, ptr %33, i32 0, i32 1
  %126 = load ptr, ptr %6, align 8, !tbaa !5
  call void @_ZNSt10shared_ptrIN7hipsycl2rt8dag_nodeEEC2ERKS3_(ptr noundef nonnull align 8 dereferenceable(16) %125, ptr noundef nonnull align 8 dereferenceable(16) %126) #6
  invoke void @"_ZNSt8functionIFvvEEC2IZN7hipsycl2rt9ocl_queue24submit_external_wait_forERKSt10shared_ptrINS4_8dag_nodeEEE3$_0vEEOT_"(ptr noundef nonnull align 8 dereferenceable(32) %32, ptr noundef nonnull align 8 dereferenceable(24) %33)
          to label %127 unwind label %169

127:                                              ; preds = %124
  invoke void @_ZN7hipsycl2rt13worker_threadclESt8functionIFvvEE(ptr noundef nonnull align 8 dereferenceable(184) %122, ptr noundef %32)
          to label %128 unwind label %173

128:                                              ; preds = %127
  call void @_ZNSt14_Function_baseD2Ev(ptr noundef nonnull align 8 dereferenceable(32) %32) #6
  call void @"_ZZN7hipsycl2rt9ocl_queue24submit_external_wait_forERKSt10shared_ptrINS0_8dag_nodeEEEN3$_0D2Ev"(ptr noundef nonnull align 8 dereferenceable(24) %33) #6
  call void @llvm.lifetime.end.p0(i64 24, ptr %33) #6
  invoke void @_ZN7hipsycl2rt12make_successEv(ptr dead_on_unwind writable sret(%"class.hipsycl::rt::result") align 8 %0)
          to label %129 unwind label %157

129:                                              ; preds = %128
  store i32 1, ptr %24, align 4
  call void @_ZNSt6vectorIN2cl5EventESaIS1_EED2Ev(ptr noundef nonnull align 8 dereferenceable(24) %26) #6
  call void @llvm.lifetime.end.p0(i64 24, ptr %26) #6
  call void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %25) #6
  call void @llvm.lifetime.end.p0(i64 8, ptr %25) #6
  br label %182

130:                                              ; preds = %100
  %131 = landingpad { ptr, i32 }
          cleanup
  %132 = extractvalue { ptr, i32 } %131, 0
  store ptr %132, ptr %11, align 8
  %133 = extractvalue { ptr, i32 } %131, 1
  store i32 %133, ptr %12, align 4
  br label %181

134:                                              ; preds = %101
  %135 = landingpad { ptr, i32 }
          cleanup
  %136 = extractvalue { ptr, i32 } %135, 0
  store ptr %136, ptr %11, align 8
  %137 = extractvalue { ptr, i32 } %135, 1
  store i32 %137, ptr %12, align 4
  %138 = load ptr, ptr %29, align 8
  %139 = icmp eq ptr %28, %138
  br i1 %139, label %144, label %140

140:                                              ; preds = %140, %134
  %141 = phi ptr [ %138, %134 ], [ %142, %140 ]
  %142 = getelementptr inbounds %"class.cl::Event", ptr %141, i64 -1
  call void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %142) #6
  %143 = icmp eq ptr %142, %28
  br i1 %143, label %144, label %140

144:                                              ; preds = %140, %134
  br label %156

145:                                              ; preds = %102
  %146 = landingpad { ptr, i32 }
          cleanup
  %147 = extractvalue { ptr, i32 } %146, 0
  store ptr %147, ptr %11, align 8
  %148 = extractvalue { ptr, i32 } %146, 1
  store i32 %148, ptr %12, align 4
  call void @_ZNSt15__new_allocatorIN2cl5EventEED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %30) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr %30) #6
  %149 = getelementptr inbounds [1 x %"class.cl::Event"], ptr %28, i32 0, i32 0
  %150 = getelementptr inbounds %"class.cl::Event", ptr %149, i64 1
  br label %151

151:                                              ; preds = %151, %145
  %152 = phi ptr [ %150, %145 ], [ %153, %151 ]
  %153 = getelementptr inbounds %"class.cl::Event", ptr %152, i64 -1
  call void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %153) #6
  %154 = icmp eq ptr %153, %149
  br i1 %154, label %155, label %151

155:                                              ; preds = %151
  br label %156

156:                                              ; preds = %155, %144
  call void @llvm.lifetime.end.p0(i64 8, ptr %28) #6
  br label %180

157:                                              ; preds = %128, %119, %116
  %158 = landingpad { ptr, i32 }
          cleanup
  %159 = extractvalue { ptr, i32 } %158, 0
  store ptr %159, ptr %11, align 8
  %160 = extractvalue { ptr, i32 } %158, 1
  store i32 %160, ptr %12, align 4
  br label %179

161:                                              ; preds = %120
  %162 = landingpad { ptr, i32 }
          cleanup
  %163 = extractvalue { ptr, i32 } %162, 0
  store ptr %163, ptr %11, align 8
  %164 = extractvalue { ptr, i32 } %162, 1
  store i32 %164, ptr %12, align 4
  call void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %31) #6
  br label %179

165:                                              ; preds = %121
  %166 = landingpad { ptr, i32 }
          cleanup
  %167 = extractvalue { ptr, i32 } %166, 0
  store ptr %167, ptr %11, align 8
  %168 = extractvalue { ptr, i32 } %166, 1
  store i32 %168, ptr %12, align 4
  br label %178

169:                                              ; preds = %124
  %170 = landingpad { ptr, i32 }
          cleanup
  %171 = extractvalue { ptr, i32 } %170, 0
  store ptr %171, ptr %11, align 8
  %172 = extractvalue { ptr, i32 } %170, 1
  store i32 %172, ptr %12, align 4
  br label %177

173:                                              ; preds = %127
  %174 = landingpad { ptr, i32 }
          cleanup
  %175 = extractvalue { ptr, i32 } %174, 0
  store ptr %175, ptr %11, align 8
  %176 = extractvalue { ptr, i32 } %174, 1
  store i32 %176, ptr %12, align 4
  call void @_ZNSt14_Function_baseD2Ev(ptr noundef nonnull align 8 dereferenceable(32) %32) #6
  br label %177

177:                                              ; preds = %173, %169
  call void @"_ZZN7hipsycl2rt9ocl_queue24submit_external_wait_forERKSt10shared_ptrINS0_8dag_nodeEEEN3$_0D2Ev"(ptr noundef nonnull align 8 dereferenceable(24) %33) #6
  br label %178

178:                                              ; preds = %177, %165
  call void @llvm.lifetime.end.p0(i64 24, ptr %33) #6
  br label %179

179:                                              ; preds = %178, %161, %157
  call void @_ZNSt6vectorIN2cl5EventESaIS1_EED2Ev(ptr noundef nonnull align 8 dereferenceable(24) %26) #6
  br label %180

180:                                              ; preds = %179, %156
  call void @llvm.lifetime.end.p0(i64 24, ptr %26) #6
  call void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %25) #6
  br label %181

181:                                              ; preds = %180, %130
  call void @llvm.lifetime.end.p0(i64 8, ptr %25) #6
  br label %183

182:                                              ; preds = %129, %56
  call void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %9) #6
  call void @llvm.lifetime.end.p0(i64 8, ptr %9) #6
  call void @llvm.lifetime.end.p0(i64 4, ptr %8) #6
  call void @llvm.lifetime.end.p0(i64 8, ptr %7) #6
  ret void

183:                                              ; preds = %181, %99
  call void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %9) #6
  br label %184

184:                                              ; preds = %183, %57
  call void @llvm.lifetime.end.p0(i64 8, ptr %9) #6
  call void @llvm.lifetime.end.p0(i64 4, ptr %8) #6
  call void @llvm.lifetime.end.p0(i64 8, ptr %7) #6
  br label %185

185:                                              ; preds = %184
  %186 = load ptr, ptr %11, align 8
  %187 = load i32, ptr %12, align 4
  %188 = insertvalue { ptr, i32 } poison, ptr %186, 0
  %189 = insertvalue { ptr, i32 } %188, i32 %187, 1
  resume { ptr, i32 } %189
}

; Function Attrs: mustprogress uwtable
declare dso_local void @_ZN2cl9UserEventC2ERKNS_7ContextEPi(ptr noundef nonnull align 8 dereferenceable(8), ptr noundef nonnull align 8 dereferenceable(8), ptr noundef) unnamed_addr #3 align 2

declare void @_ZN7hipsycl2rt13worker_threadclESt8functionIFvvEE(ptr noundef nonnull align 8 dereferenceable(184), ptr noundef) #1

; Function Attrs: inlinehint mustprogress uwtable
declare dso_local void @_ZN2cl9UserEventC2ERKS0_(ptr noundef nonnull align 8 dereferenceable(8), ptr noundef nonnull align 8 dereferenceable(8)) unnamed_addr #5 align 2

; Function Attrs: mustprogress nounwind uwtable
declare dso_local void @_ZNSt10shared_ptrIN7hipsycl2rt8dag_nodeEEC2ERKS3_(ptr noundef nonnull align 8 dereferenceable(16), ptr noundef nonnull align 8 dereferenceable(16)) unnamed_addr #2 align 2

; Function Attrs: mustprogress uwtable
declare hidden void @"_ZNSt8functionIFvvEEC2IZN7hipsycl2rt9ocl_queue24submit_external_wait_forERKSt10shared_ptrINS4_8dag_nodeEEE3$_0vEEOT_"(ptr noundef nonnull align 8 dereferenceable(32), ptr noundef nonnull align 8 dereferenceable(24)) unnamed_addr #3 align 2

; Function Attrs: mustprogress nounwind uwtable
declare dso_local void @_ZNSt14_Function_baseD2Ev(ptr noundef nonnull align 8 dereferenceable(24)) unnamed_addr #2 align 2

; Function Attrs: inlinehint mustprogress nounwind uwtable
declare hidden void @"_ZZN7hipsycl2rt9ocl_queue24submit_external_wait_forERKSt10shared_ptrINS0_8dag_nodeEEEN3$_0D2Ev"(ptr noundef nonnull align 8 dereferenceable(24)) unnamed_addr #4 align 2

; Function Attrs: mustprogress nounwind uwtable
declare dso_local void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8)) unnamed_addr #2 align 2

; Function Attrs: mustprogress nounwind uwtable
declare dso_local void @_ZNSt15__new_allocatorIN2cl5EventEED2Ev(ptr noundef nonnull align 1 dereferenceable(1)) unnamed_addr #2 align 2

attributes #0 = { nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) }
attributes #1 = { "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cmov,+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #2 = { mustprogress nounwind uwtable "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cmov,+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #3 = { mustprogress uwtable "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cmov,+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #4 = { inlinehint mustprogress nounwind uwtable "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cmov,+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #5 = { inlinehint mustprogress uwtable "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cmov,+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #6 = { nounwind }

!llvm.module.flags = !{!0, !1, !2, !3}
!llvm.ident = !{!4}

!0 = !{i32 1, !"wchar_size", i32 4}
!1 = !{i32 8, !"PIC Level", i32 2}
!2 = !{i32 7, !"PIE Level", i32 2}
!3 = !{i32 7, !"uwtable", i32 2}
!4 = !{!"clang version 19.1.0 (git@github.com:llvm/llvm-project.git a4bf6cd7cfb1a1421ba92bca9d017b49936c55e4)"}
!5 = !{!6, !6, i64 0}
!6 = !{!"any pointer", !7, i64 0}
!7 = !{!"omnipotent char", !8, i64 0}
!8 = !{!"Simple C++ TBAA"}
!9 = !{!10, !6, i64 8}
!10 = !{!"_ZTSN7hipsycl2rt9ocl_queueE", !11, i64 0, !6, i64 8, !12, i64 16, !13, i64 24, !15, i64 32, !17, i64 48, !33, i64 232, !36, i64 248, !39, i64 304, !42, i64 312, !51, i64 368, !73, i64 528}
!11 = !{!"_ZTSN7hipsycl2rt13inorder_queueE"}
!12 = !{!"long", !7, i64 0}
!13 = !{!"_ZTSN2cl12CommandQueueE", !14, i64 0}
!14 = !{!"_ZTSN2cl6detail7WrapperIP17_cl_command_queueEE", !6, i64 0}
!15 = !{!"_ZTSN7hipsycl2rt28ocl_sscp_code_object_invokerE", !16, i64 0, !6, i64 8}
!16 = !{!"_ZTSN7hipsycl2rt24sscp_code_object_invokerE"}
!17 = !{!"_ZTSN7hipsycl2rt13worker_threadE", !18, i64 0, !20, i64 8, !23, i64 16, !25, i64 64, !27, i64 104}
!18 = !{!"_ZTSSt6thread", !19, i64 0}
!19 = !{!"_ZTSNSt6thread2idE", !12, i64 0}
!20 = !{!"_ZTSSt6atomicIbE", !21, i64 0}
!21 = !{!"_ZTSSt13__atomic_baseIbE", !22, i64 0}
!22 = !{!"bool", !7, i64 0}
!23 = !{!"_ZTSSt18condition_variable", !24, i64 0}
!24 = !{!"_ZTSSt9__condvar", !7, i64 0}
!25 = !{!"_ZTSSt5mutex", !26, i64 0}
!26 = !{!"_ZTSSt12__mutex_base", !7, i64 0}
!27 = !{!"_ZTSSt5queueISt8functionIFvvEESt5dequeIS2_SaIS2_EEE", !28, i64 0}
!28 = !{!"_ZTSSt5dequeISt8functionIFvvEESaIS2_EE", !29, i64 0}
!29 = !{!"_ZTSSt11_Deque_baseISt8functionIFvvEESaIS2_EE", !30, i64 0}
!30 = !{!"_ZTSNSt11_Deque_baseISt8functionIFvvEESaIS2_EE11_Deque_implE", !31, i64 0}
!31 = !{!"_ZTSNSt11_Deque_baseISt8functionIFvvEESaIS2_EE16_Deque_impl_dataE", !6, i64 0, !12, i64 8, !32, i64 16, !32, i64 48}
!32 = !{!"_ZTSSt15_Deque_iteratorISt8functionIFvvEERS2_PS2_E", !6, i64 0, !6, i64 8, !6, i64 16, !6, i64 24}
!33 = !{!"_ZTSSt10shared_ptrIN7hipsycl2rt12kernel_cacheEE", !34, i64 0}
!34 = !{!"_ZTSSt12__shared_ptrIN7hipsycl2rt12kernel_cacheELN9__gnu_cxx12_Lock_policyE2EE", !6, i64 0, !35, i64 8}
!35 = !{!"_ZTSSt14__shared_countILN9__gnu_cxx12_Lock_policyE2EE", !6, i64 0}
!36 = !{!"_ZTSN7hipsycl2rt9ocl_queue15protected_stateE", !37, i64 0, !25, i64 16}
!37 = !{!"_ZTSSt10shared_ptrIN7hipsycl2rt14dag_node_eventEE", !38, i64 0}
!38 = !{!"_ZTSSt12__shared_ptrIN7hipsycl2rt14dag_node_eventELN9__gnu_cxx12_Lock_policyE2EE", !6, i64 0, !35, i64 8}
!39 = !{!"_ZTSN7hipsycl6common9spin_lockE", !40, i64 0}
!40 = !{!"_ZTSSt11atomic_flag", !41, i64 0}
!41 = !{!"_ZTSSt18__atomic_flag_base", !22, i64 0}
!42 = !{!"_ZTSN7hipsycl4glue3jit19cxx_argument_mapperE", !22, i64 0, !43, i64 8, !47, i64 32}
!43 = !{!"_ZTSSt6vectorIPvSaIS0_EE", !44, i64 0}
!44 = !{!"_ZTSSt12_Vector_baseIPvSaIS0_EE", !45, i64 0}
!45 = !{!"_ZTSNSt12_Vector_baseIPvSaIS0_EE12_Vector_implE", !46, i64 0}
!46 = !{!"_ZTSNSt12_Vector_baseIPvSaIS0_EE17_Vector_impl_dataE", !6, i64 0, !6, i64 8, !6, i64 16}
!47 = !{!"_ZTSSt6vectorImSaImEE", !48, i64 0}
!48 = !{!"_ZTSSt12_Vector_baseImSaImEE", !49, i64 0}
!49 = !{!"_ZTSNSt12_Vector_baseImSaImEE12_Vector_implE", !50, i64 0}
!50 = !{!"_ZTSNSt12_Vector_baseImSaImEE17_Vector_impl_dataE", !6, i64 0, !6, i64 8, !6, i64 16}
!51 = !{!"_ZTSN7hipsycl2rt20kernel_configurationE", !52, i64 0, !56, i64 24, !60, i64 48, !64, i64 72, !47, i64 96, !68, i64 120, !72, i64 144}
!52 = !{!"_ZTSSt6vectorIN7hipsycl2rt17kernel_build_flagESaIS2_EE", !53, i64 0}
!53 = !{!"_ZTSSt12_Vector_baseIN7hipsycl2rt17kernel_build_flagESaIS2_EE", !54, i64 0}
!54 = !{!"_ZTSNSt12_Vector_baseIN7hipsycl2rt17kernel_build_flagESaIS2_EE12_Vector_implE", !55, i64 0}
!55 = !{!"_ZTSNSt12_Vector_baseIN7hipsycl2rt17kernel_build_flagESaIS2_EE17_Vector_impl_dataE", !6, i64 0, !6, i64 8, !6, i64 16}
!56 = !{!"_ZTSSt6vectorISt4pairIN7hipsycl2rt19kernel_build_optionENS2_20kernel_configuration13int_or_stringEESaIS6_EE", !57, i64 0}
!57 = !{!"_ZTSSt12_Vector_baseISt4pairIN7hipsycl2rt19kernel_build_optionENS2_20kernel_configuration13int_or_stringEESaIS6_EE", !58, i64 0}
!58 = !{!"_ZTSNSt12_Vector_baseISt4pairIN7hipsycl2rt19kernel_build_optionENS2_20kernel_configuration13int_or_stringEESaIS6_EE12_Vector_implE", !59, i64 0}
!59 = !{!"_ZTSNSt12_Vector_baseISt4pairIN7hipsycl2rt19kernel_build_optionENS2_20kernel_configuration13int_or_stringEESaIS6_EE17_Vector_impl_dataE", !6, i64 0, !6, i64 8, !6, i64 16}
!60 = !{!"_ZTSSt6vectorISt4pairIimESaIS1_EE", !61, i64 0}
!61 = !{!"_ZTSSt12_Vector_baseISt4pairIimESaIS1_EE", !62, i64 0}
!62 = !{!"_ZTSNSt12_Vector_baseISt4pairIimESaIS1_EE12_Vector_implE", !63, i64 0}
!63 = !{!"_ZTSNSt12_Vector_baseISt4pairIimESaIS1_EE17_Vector_impl_dataE", !6, i64 0, !6, i64 8, !6, i64 16}
!64 = !{!"_ZTSSt6vectorIN7hipsycl4glue4sscp63__acpp_sscp_emit_param_type_annotation_fcall_specialized_configIPKNS2_24fcall_specialized_configEEESaIS7_EE", !65, i64 0}
!65 = !{!"_ZTSSt12_Vector_baseIN7hipsycl4glue4sscp63__acpp_sscp_emit_param_type_annotation_fcall_specialized_configIPKNS2_24fcall_specialized_configEEESaIS7_EE", !66, i64 0}
!66 = !{!"_ZTSNSt12_Vector_baseIN7hipsycl4glue4sscp63__acpp_sscp_emit_param_type_annotation_fcall_specialized_configIPKNS2_24fcall_specialized_configEEESaIS7_EE12_Vector_implE", !67, i64 0}
!67 = !{!"_ZTSNSt12_Vector_baseIN7hipsycl4glue4sscp63__acpp_sscp_emit_param_type_annotation_fcall_specialized_configIPKNS2_24fcall_specialized_configEEESaIS7_EE17_Vector_impl_dataE", !6, i64 0, !6, i64 8, !6, i64 16}
!68 = !{!"_ZTSSt6vectorISt4pairIiiESaIS1_EE", !69, i64 0}
!69 = !{!"_ZTSSt12_Vector_baseISt4pairIiiESaIS1_EE", !70, i64 0}
!70 = !{!"_ZTSNSt12_Vector_baseISt4pairIiiESaIS1_EE12_Vector_implE", !71, i64 0}
!71 = !{!"_ZTSNSt12_Vector_baseISt4pairIiiESaIS1_EE17_Vector_impl_dataE", !6, i64 0, !6, i64 8, !6, i64 16}
!72 = !{!"_ZTSSt5arrayImLm2EE", !7, i64 0}
!73 = !{!"_ZTSSt13unordered_mapINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEmSt4hashIS5_ESt8equal_toIS5_ESaISt4pairIKS5_mEEE", !74, i64 0}
!74 = !{!"_ZTSSt10_HashtableINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESt4pairIKS5_mESaIS8_ENSt8__detail10_Select1stESt8equal_toIS5_ESt4hashIS5_ENSA_18_Mod_range_hashingENSA_20_Default_ranged_hashENSA_20_Prime_rehash_policyENSA_17_Hashtable_traitsILb1ELb0ELb1EEEE", !6, i64 0, !12, i64 8, !75, i64 16, !12, i64 24, !76, i64 32, !6, i64 48}
!75 = !{!"_ZTSNSt8__detail15_Hash_node_baseE", !6, i64 0}
!76 = !{!"_ZTSNSt8__detail20_Prime_rehash_policyE", !77, i64 0, !12, i64 8}
!77 = !{!"float", !7, i64 0}
!78 = !{!10, !12, i64 16}
!79 = !{!80, !80, i64 0}
!80 = !{!"vtable pointer", !8, i64 0}
!81 = !{!82, !82, i64 0}
!82 = !{!"int", !7, i64 0}
!83 = !{!84, !6, i64 0}
!84 = !{!"_ZTSSt16initializer_listIN2cl5EventEE", !6, i64 0, !12, i64 8}
!85 = !{!84, !12, i64 8}


; Optimized IR (opt -passes='default<O3>'):
target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-i128:128-f80:128-n8:16:32:64-S128"
target triple = "x86_64-unknown-linux-gnu"

%"class.cl::Context" = type { %"class.cl::detail::Wrapper.51" }
%"class.cl::detail::Wrapper.51" = type { ptr }
%"class.hipsycl::rt::result" = type { %"class.std::unique_ptr" }
%"class.std::unique_ptr" = type { %"struct.std::__uniq_ptr_data" }
%"struct.std::__uniq_ptr_data" = type { %"class.std::__uniq_ptr_impl" }
%"class.std::__uniq_ptr_impl" = type { %"class.std::tuple" }
%"class.std::tuple" = type { %"struct.std::_Tuple_impl" }
%"struct.std::_Tuple_impl" = type { %"struct.std::_Head_base.57" }
%"struct.std::_Head_base.57" = type { ptr }
%"class.cl::UserEvent" = type { %"class.cl::Event" }
%"class.cl::Event" = type { %"class.cl::detail::Wrapper.61" }
%"class.cl::detail::Wrapper.61" = type { ptr }
%"class.hipsycl::rt::source_location" = type <{ %"class.std::__cxx11::basic_string", %"class.std::__cxx11::basic_string", i32, [4 x i8] }>
%"class.std::__cxx11::basic_string" = type { %"struct.std::__cxx11::basic_string<char>::_Alloc_hider", i64, %union.anon }
%"struct.std::__cxx11::basic_string<char>::_Alloc_hider" = type { ptr }
%union.anon = type { i64, [8 x i8] }
%"class.std::allocator.58" = type { i8 }
%"class.hipsycl::rt::error_info" = type <{ %"class.std::__cxx11::basic_string", %"class.hipsycl::rt::error_code", i32, [4 x i8] }>
%"class.hipsycl::rt::error_code" = type { %"class.std::__cxx11::basic_string", i8, i32 }
%"class.std::vector.74" = type { %"struct.std::_Vector_base.75" }
%"struct.std::_Vector_base.75" = type { %"struct.std::_Vector_base<cl::Event, std::allocator<cl::Event>>::_Vector_impl" }
%"struct.std::_Vector_base<cl::Event, std::allocator<cl::Event>>::_Vector_impl" = type { %"struct.std::_Vector_base<cl::Event, std::allocator<cl::Event>>::_Vector_impl_data" }
%"struct.std::_Vector_base<cl::Event, std::allocator<cl::Event>>::_Vector_impl_data" = type { ptr, ptr, ptr }
%"class.std::allocator.76" = type { i8 }
%"class.std::function.116" = type { %"class.std::_Function_base", ptr }
%"class.std::_Function_base" = type { %"union.std::_Any_data", ptr }
%"union.std::_Any_data" = type { %"union.std::_Nocopy_types" }
%"union.std::_Nocopy_types" = type { { i64, i64 } }
%class.anon.118 = type { %"class.cl::UserEvent", %"class.std::shared_ptr.92" }
%"class.std::shared_ptr.92" = type { %"class.std::__shared_ptr.93" }
%"class.std::__shared_ptr.93" = type { ptr, %"class.std::__shared_count" }
%"class.std::__shared_count" = type { ptr }

@.str = external hidden unnamed_addr constant [87 x i8], align 1
@.str.3 = external hidden unnamed_addr constant [3 x i8], align 1
@__func__._ZN7hipsycl2rt9ocl_queue24submit_external_wait_forERKSt10shared_ptrINS0_8dag_nodeEE = external hidden unnamed_addr constant [25 x i8], align 1
@.str.17 = external hidden unnamed_addr constant [45 x i8], align 1

declare i32 @__gxx_personality_v0(...)

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

declare void @_ZNK7hipsycl2rt20ocl_hardware_context14get_cl_contextEv(ptr dead_on_unwind writable sret(%"class.cl::Context") align 8, ptr noundef nonnull align 8 dereferenceable(89)) local_unnamed_addr #1

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

; Function Attrs: mustprogress nounwind uwtable
declare void @_ZNSaIcEC2Ev(ptr noundef nonnull align 1 dereferenceable(1)) unnamed_addr #2 align 2

; Function Attrs: mustprogress uwtable
declare dso_local void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_(ptr noundef nonnull align 8 dereferenceable(32), ptr noundef, ptr noundef nonnull align 1 dereferenceable(1)) unnamed_addr #3 align 2

; Function Attrs: mustprogress uwtable
declare dso_local void @_ZN7hipsycl2rt15source_locationC2ERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEES9_i(ptr noundef nonnull align 8 dereferenceable(68), ptr noundef nonnull align 8 dereferenceable(32), ptr noundef nonnull align 8 dereferenceable(32), i32 noundef) unnamed_addr #3 align 2

; Function Attrs: mustprogress uwtable
declare dso_local void @_ZN7hipsycl2rt10error_codeC2ERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEi(ptr noundef nonnull align 8 dereferenceable(40), ptr noundef nonnull align 8 dereferenceable(32), i32 noundef) unnamed_addr #3 align 2

; Function Attrs: mustprogress uwtable
declare dso_local void @_ZN7hipsycl2rt10error_infoC2ERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEENS0_10error_codeENS0_10error_typeE(ptr noundef nonnull align 8 dereferenceable(76), ptr noundef nonnull align 8 dereferenceable(32), ptr noundef, i32 noundef) unnamed_addr #3 align 2

; Function Attrs: inlinehint mustprogress nounwind uwtable
declare dso_local void @_ZN7hipsycl2rt10error_infoD2Ev(ptr noundef nonnull align 8 dereferenceable(76)) unnamed_addr #4 align 2

; Function Attrs: inlinehint mustprogress nounwind uwtable
declare dso_local void @_ZN7hipsycl2rt10error_codeD2Ev(ptr noundef nonnull align 8 dereferenceable(40)) unnamed_addr #4 align 2

; Function Attrs: mustprogress nounwind uwtable
declare void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32)) unnamed_addr #2 align 2

; Function Attrs: mustprogress nounwind uwtable
declare dso_local void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1)) unnamed_addr #2 align 2

; Function Attrs: inlinehint mustprogress nounwind uwtable
declare dso_local void @_ZN7hipsycl2rt15source_locationD2Ev(ptr noundef nonnull align 8 dereferenceable(68)) unnamed_addr #4 align 2

; Function Attrs: mustprogress nounwind uwtable
declare dso_local void @_ZN2cl6detail7WrapperIP11_cl_contextED2Ev(ptr noundef nonnull align 8 dereferenceable(8)) unnamed_addr #2 align 2

; Function Attrs: mustprogress uwtable
declare dso_local void @_ZN2cl5EventC2Ev(ptr noundef nonnull align 8 dereferenceable(8)) unnamed_addr #3 align 2

; Function Attrs: mustprogress uwtable
declare dso_local noundef i32 @_ZNK2cl12CommandQueue26enqueueBarrierWithWaitListEPKSt6vectorINS_5EventESaIS2_EEPS2_(ptr noundef nonnull align 8 dereferenceable(8), ptr noundef, ptr noundef) local_unnamed_addr #3 align 2

; Function Attrs: mustprogress uwtable
declare dso_local void @_ZN7hipsycl2rt9ocl_queue21register_submitted_opEN2cl5EventE(ptr noundef nonnull align 8 dereferenceable(584), ptr noundef) local_unnamed_addr #3 align 2

; Function Attrs: inlinehint mustprogress uwtable
declare dso_local void @_ZN2cl5EventC2ERKS0_(ptr noundef nonnull align 8 dereferenceable(8), ptr noundef nonnull align 8 dereferenceable(8)) unnamed_addr #5 align 2

; Function Attrs: mustprogress nounwind uwtable
declare dso_local void @_ZNSt6vectorIN2cl5EventESaIS1_EED2Ev(ptr noundef nonnull align 8 dereferenceable(24)) unnamed_addr #2 align 2

; Function Attrs: inlinehint mustprogress uwtable
declare dso_local void @_ZN7hipsycl2rt10make_errorERKNS0_15source_locationERKNS0_10error_infoE(ptr dead_on_unwind noalias writable sret(%"class.hipsycl::rt::result") align 8, ptr noundef nonnull align 8 dereferenceable(68), ptr noundef nonnull align 8 dereferenceable(76)) local_unnamed_addr #5

; Function Attrs: inlinehint mustprogress nounwind uwtable
declare dso_local void @_ZN7hipsycl2rt12make_successEv(ptr dead_on_unwind noalias writable sret(%"class.hipsycl::rt::result") align 8) local_unnamed_addr #4

; Function Attrs: mustprogress nounwind uwtable
declare dso_local void @_ZNSaIN2cl5EventEEC2Ev(ptr noundef nonnull align 1 dereferenceable(1)) unnamed_addr #2 align 2

; Function Attrs: mustprogress uwtable
declare dso_local void @_ZNSt6vectorIN2cl5EventESaIS1_EEC2ESt16initializer_listIS1_ERKS2_(ptr noundef nonnull align 8 dereferenceable(24), ptr, i64, ptr noundef nonnull align 1 dereferenceable(1)) unnamed_addr #3 align 2

; Function Attrs: mustprogress uwtable
define dso_local void @_ZN7hipsycl2rt9ocl_queue24submit_external_wait_forERKSt10shared_ptrINS0_8dag_nodeEE(ptr dead_on_unwind noalias writable sret(%"class.hipsycl::rt::result") align 8 %0, ptr noundef nonnull align 8 dereferenceable(584) %1, ptr noundef nonnull align 8 dereferenceable(16) %2) unnamed_addr #3 align 2 personality ptr @__gxx_personality_v0 {
  %4 = alloca i32, align 4
  %5 = alloca %"class.cl::UserEvent", align 8
  %6 = alloca %"class.cl::Context", align 8
  %7 = alloca %"class.hipsycl::rt::source_location", align 8
  %8 = alloca %"class.std::__cxx11::basic_string", align 8
  %9 = alloca %"class.std::allocator.58", align 1
  %10 = alloca %"class.std::__cxx11::basic_string", align 8
  %11 = alloca %"class.std::allocator.58", align 1
  %12 = alloca %"class.hipsycl::rt::error_info", align 8
  %13 = alloca %"class.std::__cxx11::basic_string", align 8
  %14 = alloca %"class.std::allocator.58", align 1
  %15 = alloca %"class.hipsycl::rt::error_code", align 8
  %16 = alloca %"class.std::__cxx11::basic_string", align 8
  %17 = alloca %"class.std::allocator.58", align 1
  %18 = alloca %"class.cl::Event", align 8
  %19 = alloca %"class.std::vector.74", align 8
  %20 = alloca [1 x %"class.cl::Event"], align 8
  %21 = alloca %"class.std::allocator.76", align 1
  %22 = alloca %"class.cl::Event", align 8
  %23 = alloca %"class.std::function.116", align 8
  %24 = alloca %class.anon.118, align 8
  %25 = getelementptr inbounds i8, ptr %1, i64 8
  %26 = load ptr, ptr %25, align 8, !tbaa !5
  %27 = getelementptr inbounds i8, ptr %1, i64 16
  %28 = load i64, ptr %27, align 8, !tbaa !77
  %29 = load ptr, ptr %26, align 8, !tbaa !78
  %30 = getelementptr inbounds i8, ptr %29, i64 16
  %31 = load ptr, ptr %30, align 8
  %32 = tail call noundef ptr %31(ptr noundef nonnull align 8 dereferenceable(60) %26, i64 noundef %28)
  call void @llvm.lifetime.start.p0(i64 4, ptr nonnull %4) #6
  call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %5) #6
  call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %6) #6
  call void @_ZNK7hipsycl2rt20ocl_hardware_context14get_cl_contextEv(ptr dead_on_unwind nonnull writable sret(%"class.cl::Context") align 8 %6, ptr noundef nonnull align 8 dereferenceable(89) %32)
  invoke void @_ZN2cl9UserEventC2ERKNS_7ContextEPi(ptr noundef nonnull align 8 dereferenceable(8) %5, ptr noundef nonnull align 8 dereferenceable(8) %6, ptr noundef nonnull %4)
          to label %33 unwind label %45

33:                                               ; preds = %3
  call void @_ZN2cl6detail7WrapperIP11_cl_contextED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %6) #6
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %6) #6
  %34 = load i32, ptr %4, align 4, !tbaa !80
  %.not = icmp eq i32 %34, 0
  br i1 %.not, label %70, label %35

35:                                               ; preds = %33
  call void @llvm.lifetime.start.p0(i64 72, ptr nonnull %7) #6
  call void @llvm.lifetime.start.p0(i64 32, ptr nonnull %8) #6
  call void @llvm.lifetime.start.p0(i64 1, ptr nonnull %9) #6
  call void @_ZNSaIcEC2Ev(ptr noundef nonnull align 1 dereferenceable(1) %9) #6
  invoke void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_(ptr noundef nonnull align 8 dereferenceable(32) %8, ptr noundef nonnull @__func__._ZN7hipsycl2rt9ocl_queue24submit_external_wait_forERKSt10shared_ptrINS0_8dag_nodeEE, ptr noundef nonnull align 1 dereferenceable(1) %9)
          to label %36 unwind label %47

36:                                               ; preds = %35
  call void @llvm.lifetime.start.p0(i64 32, ptr nonnull %10) #6
  call void @llvm.lifetime.start.p0(i64 1, ptr nonnull %11) #6
  call void @_ZNSaIcEC2Ev(ptr noundef nonnull align 1 dereferenceable(1) %11) #6
  invoke void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_(ptr noundef nonnull align 8 dereferenceable(32) %10, ptr noundef nonnull @.str, ptr noundef nonnull align 1 dereferenceable(1) %11)
          to label %37 unwind label %49

37:                                               ; preds = %36
  invoke void @_ZN7hipsycl2rt15source_locationC2ERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEES9_i(ptr noundef nonnull align 8 dereferenceable(68) %7, ptr noundef nonnull align 8 dereferenceable(32) %8, ptr noundef nonnull align 8 dereferenceable(32) %10, i32 noundef 375)
          to label %38 unwind label %51

38:                                               ; preds = %37
  call void @llvm.lifetime.start.p0(i64 80, ptr nonnull %12) #6
  call void @llvm.lifetime.start.p0(i64 32, ptr nonnull %13) #6
  call void @llvm.lifetime.start.p0(i64 1, ptr nonnull %14) #6
  call void @_ZNSaIcEC2Ev(ptr noundef nonnull align 1 dereferenceable(1) %14) #6
  invoke void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_(ptr noundef nonnull align 8 dereferenceable(32) %13, ptr noundef nonnull @.str.17, ptr noundef nonnull align 1 dereferenceable(1) %14)
          to label %39 unwind label %53

39:                                               ; preds = %38
  call void @llvm.lifetime.start.p0(i64 32, ptr nonnull %16) #6
  call void @llvm.lifetime.start.p0(i64 1, ptr nonnull %17) #6
  call void @_ZNSaIcEC2Ev(ptr noundef nonnull align 1 dereferenceable(1) %17) #6
  invoke void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_(ptr noundef nonnull align 8 dereferenceable(32) %16, ptr noundef nonnull @.str.3, ptr noundef nonnull align 1 dereferenceable(1) %17)
          to label %40 unwind label %55

40:                                               ; preds = %39
  %41 = load i32, ptr %4, align 4, !tbaa !80
  invoke void @_ZN7hipsycl2rt10error_codeC2ERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEi(ptr noundef nonnull align 8 dereferenceable(40) %15, ptr noundef nonnull align 8 dereferenceable(32) %16, i32 noundef %41)
          to label %42 unwind label %57

42:                                               ; preds = %40
  invoke void @_ZN7hipsycl2rt10error_infoC2ERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEENS0_10error_codeENS0_10error_typeE(ptr noundef nonnull align 8 dereferenceable(76) %12, ptr noundef nonnull align 8 dereferenceable(32) %13, ptr noundef nonnull %15, i32 noundef 1)
          to label %43 unwind label %59

43:                                               ; preds = %42
  invoke void @_ZN7hipsycl2rt10make_errorERKNS0_15source_locationERKNS0_10error_infoE(ptr dead_on_unwind writable sret(%"class.hipsycl::rt::result") align 8 %0, ptr noundef nonnull align 8 dereferenceable(68) %7, ptr noundef nonnull align 8 dereferenceable(76) %12)
          to label %44 unwind label %61

44:                                               ; preds = %43
  call void @_ZN7hipsycl2rt10error_infoD2Ev(ptr noundef nonnull align 8 dereferenceable(76) %12) #6
  call void @_ZN7hipsycl2rt10error_codeD2Ev(ptr noundef nonnull align 8 dereferenceable(40) %15) #6
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %16) #6
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %17) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr nonnull %17) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %16) #6
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %13) #6
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %14) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr nonnull %14) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %13) #6
  call void @llvm.lifetime.end.p0(i64 80, ptr nonnull %12) #6
  call void @_ZN7hipsycl2rt15source_locationD2Ev(ptr noundef nonnull align 8 dereferenceable(68) %7) #6
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %10) #6
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %11) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr nonnull %11) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %10) #6
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %8) #6
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %9) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr nonnull %9) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %8) #6
  call void @llvm.lifetime.end.p0(i64 72, ptr nonnull %7) #6
  br label %104

45:                                               ; preds = %3
  %46 = landingpad { ptr, i32 }
          cleanup
  call void @_ZN2cl6detail7WrapperIP11_cl_contextED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %6) #6
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %6) #6
  br label %106

47:                                               ; preds = %35
  %48 = landingpad { ptr, i32 }
          cleanup
  br label %69

49:                                               ; preds = %36
  %50 = landingpad { ptr, i32 }
          cleanup
  br label %68

51:                                               ; preds = %37
  %52 = landingpad { ptr, i32 }
          cleanup
  br label %67

53:                                               ; preds = %38
  %54 = landingpad { ptr, i32 }
          cleanup
  br label %66

55:                                               ; preds = %39
  %56 = landingpad { ptr, i32 }
          cleanup
  br label %65

57:                                               ; preds = %40
  %58 = landingpad { ptr, i32 }
          cleanup
  br label %64

59:                                               ; preds = %42
  %60 = landingpad { ptr, i32 }
          cleanup
  br label %63

61:                                               ; preds = %43
  %62 = landingpad { ptr, i32 }
          cleanup
  call void @_ZN7hipsycl2rt10error_infoD2Ev(ptr noundef nonnull align 8 dereferenceable(76) %12) #6
  br label %63

63:                                               ; preds = %61, %59
  %.pn28 = phi { ptr, i32 } [ %62, %61 ], [ %60, %59 ]
  call void @_ZN7hipsycl2rt10error_codeD2Ev(ptr noundef nonnull align 8 dereferenceable(40) %15) #6
  br label %64

64:                                               ; preds = %63, %57
  %.pn28.pn = phi { ptr, i32 } [ %.pn28, %63 ], [ %58, %57 ]
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %16) #6
  br label %65

65:                                               ; preds = %64, %55
  %.pn28.pn.pn = phi { ptr, i32 } [ %.pn28.pn, %64 ], [ %56, %55 ]
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %17) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr nonnull %17) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %16) #6
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %13) #6
  br label %66

66:                                               ; preds = %65, %53
  %.pn28.pn.pn.pn = phi { ptr, i32 } [ %.pn28.pn.pn, %65 ], [ %54, %53 ]
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %14) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr nonnull %14) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %13) #6
  call void @llvm.lifetime.end.p0(i64 80, ptr nonnull %12) #6
  call void @_ZN7hipsycl2rt15source_locationD2Ev(ptr noundef nonnull align 8 dereferenceable(68) %7) #6
  br label %67

67:                                               ; preds = %66, %51
  %.pn28.pn.pn.pn.pn = phi { ptr, i32 } [ %.pn28.pn.pn.pn, %66 ], [ %52, %51 ]
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %10) #6
  br label %68

68:                                               ; preds = %67, %49
  %.pn28.pn.pn.pn.pn.pn = phi { ptr, i32 } [ %.pn28.pn.pn.pn.pn, %67 ], [ %50, %49 ]
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %11) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr nonnull %11) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %10) #6
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %8) #6
  br label %69

69:                                               ; preds = %68, %47
  %.pn28.pn.pn.pn.pn.pn.pn = phi { ptr, i32 } [ %.pn28.pn.pn.pn.pn.pn, %68 ], [ %48, %47 ]
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %9) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr nonnull %9) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %8) #6
  call void @llvm.lifetime.end.p0(i64 72, ptr nonnull %7) #6
  br label %105

70:                                               ; preds = %33
  call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %18) #6
  invoke void @_ZN2cl5EventC2Ev(ptr noundef nonnull align 8 dereferenceable(8) %18)
          to label %71 unwind label %84

71:                                               ; preds = %70
  call void @llvm.lifetime.start.p0(i64 24, ptr nonnull %19) #6
  call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %20) #6
  invoke void @_ZN2cl5EventC2ERKS0_(ptr noundef nonnull align 8 dereferenceable(8) %20, ptr noundef nonnull align 8 dereferenceable(8) %5)
          to label %72 unwind label %86

72:                                               ; preds = %71
  call void @llvm.lifetime.start.p0(i64 1, ptr nonnull %21) #6
  call void @_ZNSaIN2cl5EventEEC2Ev(ptr noundef nonnull align 1 dereferenceable(1) %21) #6
  invoke void @_ZNSt6vectorIN2cl5EventESaIS1_EEC2ESt16initializer_listIS1_ERKS2_(ptr noundef nonnull align 8 dereferenceable(24) %19, ptr nonnull %20, i64 1, ptr noundef nonnull align 1 dereferenceable(1) %21)
          to label %73 unwind label %.loopexit.loopexit

73:                                               ; preds = %72
  call void @_ZNSt15__new_allocatorIN2cl5EventEED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %21) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr nonnull %21) #6
  call void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %20) #6
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %20) #6
  %74 = getelementptr inbounds i8, ptr %1, i64 24
  %75 = invoke noundef i32 @_ZNK2cl12CommandQueue26enqueueBarrierWithWaitListEPKSt6vectorINS_5EventESaIS2_EEPS2_(ptr noundef nonnull align 8 dereferenceable(8) %74, ptr noundef nonnull %19, ptr noundef nonnull %18)
          to label %76 unwind label %89

76:                                               ; preds = %73
  invoke void @_ZN2cl5EventC2ERKS0_(ptr noundef nonnull align 8 dereferenceable(8) %22, ptr noundef nonnull align 8 dereferenceable(8) %18)
          to label %77 unwind label %89

77:                                               ; preds = %76
  invoke void @_ZN7hipsycl2rt9ocl_queue21register_submitted_opEN2cl5EventE(ptr noundef nonnull align 8 dereferenceable(584) %1, ptr noundef nonnull %22)
          to label %78 unwind label %91

78:                                               ; preds = %77
  call void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %22) #6
  %79 = getelementptr inbounds i8, ptr %1, i64 48
  call void @llvm.lifetime.start.p0(i64 24, ptr nonnull %24) #6
  invoke void @_ZN2cl9UserEventC2ERKS0_(ptr noundef nonnull align 8 dereferenceable(8) %24, ptr noundef nonnull align 8 dereferenceable(8) %5)
          to label %80 unwind label %93

80:                                               ; preds = %78
  %81 = getelementptr inbounds i8, ptr %24, i64 8
  call void @_ZNSt10shared_ptrIN7hipsycl2rt8dag_nodeEEC2ERKS3_(ptr noundef nonnull align 8 dereferenceable(16) %81, ptr noundef nonnull align 8 dereferenceable(16) %2) #6
  invoke void @"_ZNSt8functionIFvvEEC2IZN7hipsycl2rt9ocl_queue24submit_external_wait_forERKSt10shared_ptrINS4_8dag_nodeEEE3$_0vEEOT_"(ptr noundef nonnull align 8 dereferenceable(32) %23, ptr noundef nonnull align 8 dereferenceable(24) %24)
          to label %82 unwind label %95

82:                                               ; preds = %80
  invoke void @_ZN7hipsycl2rt13worker_threadclESt8functionIFvvEE(ptr noundef nonnull align 8 dereferenceable(184) %79, ptr noundef nonnull %23)
          to label %83 unwind label %97

83:                                               ; preds = %82
  call void @_ZNSt14_Function_baseD2Ev(ptr noundef nonnull align 8 dereferenceable(32) %23) #6
  call void @"_ZZN7hipsycl2rt9ocl_queue24submit_external_wait_forERKSt10shared_ptrINS0_8dag_nodeEEEN3$_0D2Ev"(ptr noundef nonnull align 8 dereferenceable(24) %24) #6
  call void @llvm.lifetime.end.p0(i64 24, ptr nonnull %24) #6
  call void @_ZN7hipsycl2rt12make_successEv(ptr dead_on_unwind writable sret(%"class.hipsycl::rt::result") align 8 %0)
  call void @_ZNSt6vectorIN2cl5EventESaIS1_EED2Ev(ptr noundef nonnull align 8 dereferenceable(24) %19) #6
  call void @llvm.lifetime.end.p0(i64 24, ptr nonnull %19) #6
  call void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %18) #6
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %18) #6
  br label %104

84:                                               ; preds = %70
  %85 = landingpad { ptr, i32 }
          cleanup
  br label %103

86:                                               ; preds = %71
  %87 = landingpad { ptr, i32 }
          cleanup
  br label %.loopexit

.loopexit.loopexit:                               ; preds = %72
  %88 = landingpad { ptr, i32 }
          cleanup
  call void @_ZNSt15__new_allocatorIN2cl5EventEED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %21) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr nonnull %21) #6
  call void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %20) #6
  br label %.loopexit

.loopexit:                                        ; preds = %.loopexit.loopexit, %86
  %.pn = phi { ptr, i32 } [ %87, %86 ], [ %88, %.loopexit.loopexit ]
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %20) #6
  br label %102

89:                                               ; preds = %76, %73
  %90 = landingpad { ptr, i32 }
          cleanup
  br label %101

91:                                               ; preds = %77
  %92 = landingpad { ptr, i32 }
          cleanup
  call void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %22) #6
  br label %101

93:                                               ; preds = %78
  %94 = landingpad { ptr, i32 }
          cleanup
  br label %100

95:                                               ; preds = %80
  %96 = landingpad { ptr, i32 }
          cleanup
  br label %99

97:                                               ; preds = %82
  %98 = landingpad { ptr, i32 }
          cleanup
  call void @_ZNSt14_Function_baseD2Ev(ptr noundef nonnull align 8 dereferenceable(32) %23) #6
  br label %99

99:                                               ; preds = %97, %95
  %.pn22 = phi { ptr, i32 } [ %98, %97 ], [ %96, %95 ]
  call void @"_ZZN7hipsycl2rt9ocl_queue24submit_external_wait_forERKSt10shared_ptrINS0_8dag_nodeEEEN3$_0D2Ev"(ptr noundef nonnull align 8 dereferenceable(24) %24) #6
  br label %100

100:                                              ; preds = %99, %93
  %.pn22.pn = phi { ptr, i32 } [ %.pn22, %99 ], [ %94, %93 ]
  call void @llvm.lifetime.end.p0(i64 24, ptr nonnull %24) #6
  br label %101

101:                                              ; preds = %100, %91, %89
  %.pn22.pn.pn = phi { ptr, i32 } [ %.pn22.pn, %100 ], [ %92, %91 ], [ %90, %89 ]
  call void @_ZNSt6vectorIN2cl5EventESaIS1_EED2Ev(ptr noundef nonnull align 8 dereferenceable(24) %19) #6
  br label %102

102:                                              ; preds = %101, %.loopexit
  %.pn22.pn.pn.pn = phi { ptr, i32 } [ %.pn22.pn.pn, %101 ], [ %.pn, %.loopexit ]
  call void @llvm.lifetime.end.p0(i64 24, ptr nonnull %19) #6
  call void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %18) #6
  br label %103

103:                                              ; preds = %102, %84
  %.pn22.pn.pn.pn.pn = phi { ptr, i32 } [ %.pn22.pn.pn.pn, %102 ], [ %85, %84 ]
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %18) #6
  br label %105

104:                                              ; preds = %83, %44
  call void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %5) #6
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %5) #6
  call void @llvm.lifetime.end.p0(i64 4, ptr nonnull %4) #6
  ret void

105:                                              ; preds = %103, %69
  %.pn28.pn.pn.pn.pn.pn.pn.pn = phi { ptr, i32 } [ %.pn28.pn.pn.pn.pn.pn.pn, %69 ], [ %.pn22.pn.pn.pn.pn, %103 ]
  call void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %5) #6
  br label %106

106:                                              ; preds = %105, %45
  %.pn28.pn.pn.pn.pn.pn.pn.pn.pn = phi { ptr, i32 } [ %.pn28.pn.pn.pn.pn.pn.pn.pn, %105 ], [ %46, %45 ]
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %5) #6
  call void @llvm.lifetime.end.p0(i64 4, ptr nonnull %4) #6
  resume { ptr, i32 } %.pn28.pn.pn.pn.pn.pn.pn.pn.pn
}

; Function Attrs: mustprogress uwtable
declare dso_local void @_ZN2cl9UserEventC2ERKNS_7ContextEPi(ptr noundef nonnull align 8 dereferenceable(8), ptr noundef nonnull align 8 dereferenceable(8), ptr noundef) unnamed_addr #3 align 2

declare void @_ZN7hipsycl2rt13worker_threadclESt8functionIFvvEE(ptr noundef nonnull align 8 dereferenceable(184), ptr noundef) local_unnamed_addr #1

; Function Attrs: inlinehint mustprogress uwtable
declare dso_local void @_ZN2cl9UserEventC2ERKS0_(ptr noundef nonnull align 8 dereferenceable(8), ptr noundef nonnull align 8 dereferenceable(8)) unnamed_addr #5 align 2

; Function Attrs: mustprogress nounwind uwtable
declare dso_local void @_ZNSt10shared_ptrIN7hipsycl2rt8dag_nodeEEC2ERKS3_(ptr noundef nonnull align 8 dereferenceable(16), ptr noundef nonnull align 8 dereferenceable(16)) unnamed_addr #2 align 2

; Function Attrs: mustprogress uwtable
declare hidden void @"_ZNSt8functionIFvvEEC2IZN7hipsycl2rt9ocl_queue24submit_external_wait_forERKSt10shared_ptrINS4_8dag_nodeEEE3$_0vEEOT_"(ptr noundef nonnull align 8 dereferenceable(32), ptr noundef nonnull align 8 dereferenceable(24)) unnamed_addr #3 align 2

; Function Attrs: mustprogress nounwind uwtable
declare dso_local void @_ZNSt14_Function_baseD2Ev(ptr noundef nonnull align 8 dereferenceable(24)) unnamed_addr #2 align 2

; Function Attrs: inlinehint mustprogress nounwind uwtable
declare hidden void @"_ZZN7hipsycl2rt9ocl_queue24submit_external_wait_forERKSt10shared_ptrINS0_8dag_nodeEEEN3$_0D2Ev"(ptr noundef nonnull align 8 dereferenceable(24)) unnamed_addr #4 align 2

; Function Attrs: mustprogress nounwind uwtable
declare dso_local void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8)) unnamed_addr #2 align 2

; Function Attrs: mustprogress nounwind uwtable
declare dso_local void @_ZNSt15__new_allocatorIN2cl5EventEED2Ev(ptr noundef nonnull align 1 dereferenceable(1)) unnamed_addr #2 align 2

attributes #0 = { mustprogress nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) }
attributes #1 = { "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cmov,+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #2 = { mustprogress nounwind uwtable "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cmov,+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #3 = { mustprogress uwtable "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cmov,+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #4 = { inlinehint mustprogress nounwind uwtable "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cmov,+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #5 = { inlinehint mustprogress uwtable "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cmov,+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #6 = { nounwind }

!llvm.module.flags = !{!0, !1, !2, !3}
!llvm.ident = !{!4}

!0 = !{i32 1, !"wchar_size", i32 4}
!1 = !{i32 8, !"PIC Level", i32 2}
!2 = !{i32 7, !"PIE Level", i32 2}
!3 = !{i32 7, !"uwtable", i32 2}
!4 = !{!"clang version 19.1.0 (git@github.com:llvm/llvm-project.git a4bf6cd7cfb1a1421ba92bca9d017b49936c55e4)"}
!5 = !{!6, !8, i64 8}
!6 = !{!"_ZTSN7hipsycl2rt9ocl_queueE", !7, i64 0, !8, i64 8, !11, i64 16, !12, i64 24, !14, i64 32, !16, i64 48, !32, i64 232, !35, i64 248, !38, i64 304, !41, i64 312, !50, i64 368, !72, i64 528}
!7 = !{!"_ZTSN7hipsycl2rt13inorder_queueE"}
!8 = !{!"any pointer", !9, i64 0}
!9 = !{!"omnipotent char", !10, i64 0}
!10 = !{!"Simple C++ TBAA"}
!11 = !{!"long", !9, i64 0}
!12 = !{!"_ZTSN2cl12CommandQueueE", !13, i64 0}
!13 = !{!"_ZTSN2cl6detail7WrapperIP17_cl_command_queueEE", !8, i64 0}
!14 = !{!"_ZTSN7hipsycl2rt28ocl_sscp_code_object_invokerE", !15, i64 0, !8, i64 8}
!15 = !{!"_ZTSN7hipsycl2rt24sscp_code_object_invokerE"}
!16 = !{!"_ZTSN7hipsycl2rt13worker_threadE", !17, i64 0, !19, i64 8, !22, i64 16, !24, i64 64, !26, i64 104}
!17 = !{!"_ZTSSt6thread", !18, i64 0}
!18 = !{!"_ZTSNSt6thread2idE", !11, i64 0}
!19 = !{!"_ZTSSt6atomicIbE", !20, i64 0}
!20 = !{!"_ZTSSt13__atomic_baseIbE", !21, i64 0}
!21 = !{!"bool", !9, i64 0}
!22 = !{!"_ZTSSt18condition_variable", !23, i64 0}
!23 = !{!"_ZTSSt9__condvar", !9, i64 0}
!24 = !{!"_ZTSSt5mutex", !25, i64 0}
!25 = !{!"_ZTSSt12__mutex_base", !9, i64 0}
!26 = !{!"_ZTSSt5queueISt8functionIFvvEESt5dequeIS2_SaIS2_EEE", !27, i64 0}
!27 = !{!"_ZTSSt5dequeISt8functionIFvvEESaIS2_EE", !28, i64 0}
!28 = !{!"_ZTSSt11_Deque_baseISt8functionIFvvEESaIS2_EE", !29, i64 0}
!29 = !{!"_ZTSNSt11_Deque_baseISt8functionIFvvEESaIS2_EE11_Deque_implE", !30, i64 0}
!30 = !{!"_ZTSNSt11_Deque_baseISt8functionIFvvEESaIS2_EE16_Deque_impl_dataE", !8, i64 0, !11, i64 8, !31, i64 16, !31, i64 48}
!31 = !{!"_ZTSSt15_Deque_iteratorISt8functionIFvvEERS2_PS2_E", !8, i64 0, !8, i64 8, !8, i64 16, !8, i64 24}
!32 = !{!"_ZTSSt10shared_ptrIN7hipsycl2rt12kernel_cacheEE", !33, i64 0}
!33 = !{!"_ZTSSt12__shared_ptrIN7hipsycl2rt12kernel_cacheELN9__gnu_cxx12_Lock_policyE2EE", !8, i64 0, !34, i64 8}
!34 = !{!"_ZTSSt14__shared_countILN9__gnu_cxx12_Lock_policyE2EE", !8, i64 0}
!35 = !{!"_ZTSN7hipsycl2rt9ocl_queue15protected_stateE", !36, i64 0, !24, i64 16}
!36 = !{!"_ZTSSt10shared_ptrIN7hipsycl2rt14dag_node_eventEE", !37, i64 0}
!37 = !{!"_ZTSSt12__shared_ptrIN7hipsycl2rt14dag_node_eventELN9__gnu_cxx12_Lock_policyE2EE", !8, i64 0, !34, i64 8}
!38 = !{!"_ZTSN7hipsycl6common9spin_lockE", !39, i64 0}
!39 = !{!"_ZTSSt11atomic_flag", !40, i64 0}
!40 = !{!"_ZTSSt18__atomic_flag_base", !21, i64 0}
!41 = !{!"_ZTSN7hipsycl4glue3jit19cxx_argument_mapperE", !21, i64 0, !42, i64 8, !46, i64 32}
!42 = !{!"_ZTSSt6vectorIPvSaIS0_EE", !43, i64 0}
!43 = !{!"_ZTSSt12_Vector_baseIPvSaIS0_EE", !44, i64 0}
!44 = !{!"_ZTSNSt12_Vector_baseIPvSaIS0_EE12_Vector_implE", !45, i64 0}
!45 = !{!"_ZTSNSt12_Vector_baseIPvSaIS0_EE17_Vector_impl_dataE", !8, i64 0, !8, i64 8, !8, i64 16}
!46 = !{!"_ZTSSt6vectorImSaImEE", !47, i64 0}
!47 = !{!"_ZTSSt12_Vector_baseImSaImEE", !48, i64 0}
!48 = !{!"_ZTSNSt12_Vector_baseImSaImEE12_Vector_implE", !49, i64 0}
!49 = !{!"_ZTSNSt12_Vector_baseImSaImEE17_Vector_impl_dataE", !8, i64 0, !8, i64 8, !8, i64 16}
!50 = !{!"_ZTSN7hipsycl2rt20kernel_configurationE", !51, i64 0, !55, i64 24, !59, i64 48, !63, i64 72, !46, i64 96, !67, i64 120, !71, i64 144}
!51 = !{!"_ZTSSt6vectorIN7hipsycl2rt17kernel_build_flagESaIS2_EE", !52, i64 0}
!52 = !{!"_ZTSSt12_Vector_baseIN7hipsycl2rt17kernel_build_flagESaIS2_EE", !53, i64 0}
!53 = !{!"_ZTSNSt12_Vector_baseIN7hipsycl2rt17kernel_build_flagESaIS2_EE12_Vector_implE", !54, i64 0}
!54 = !{!"_ZTSNSt12_Vector_baseIN7hipsycl2rt17kernel_build_flagESaIS2_EE17_Vector_impl_dataE", !8, i64 0, !8, i64 8, !8, i64 16}
!55 = !{!"_ZTSSt6vectorISt4pairIN7hipsycl2rt19kernel_build_optionENS2_20kernel_configuration13int_or_stringEESaIS6_EE", !56, i64 0}
!56 = !{!"_ZTSSt12_Vector_baseISt4pairIN7hipsycl2rt19kernel_build_optionENS2_20kernel_configuration13int_or_stringEESaIS6_EE", !57, i64 0}
!57 = !{!"_ZTSNSt12_Vector_baseISt4pairIN7hipsycl2rt19kernel_build_optionENS2_20kernel_configuration13int_or_stringEESaIS6_EE12_Vector_implE", !58, i64 0}
!58 = !{!"_ZTSNSt12_Vector_baseISt4pairIN7hipsycl2rt19kernel_build_optionENS2_20kernel_configuration13int_or_stringEESaIS6_EE17_Vector_impl_dataE", !8, i64 0, !8, i64 8, !8, i64 16}
!59 = !{!"_ZTSSt6vectorISt4pairIimESaIS1_EE", !60, i64 0}
!60 = !{!"_ZTSSt12_Vector_baseISt4pairIimESaIS1_EE", !61, i64 0}
!61 = !{!"_ZTSNSt12_Vector_baseISt4pairIimESaIS1_EE12_Vector_implE", !62, i64 0}
!62 = !{!"_ZTSNSt12_Vector_baseISt4pairIimESaIS1_EE17_Vector_impl_dataE", !8, i64 0, !8, i64 8, !8, i64 16}
!63 = !{!"_ZTSSt6vectorIN7hipsycl4glue4sscp63__acpp_sscp_emit_param_type_annotation_fcall_specialized_configIPKNS2_24fcall_specialized_configEEESaIS7_EE", !64, i64 0}
!64 = !{!"_ZTSSt12_Vector_baseIN7hipsycl4glue4sscp63__acpp_sscp_emit_param_type_annotation_fcall_specialized_configIPKNS2_24fcall_specialized_configEEESaIS7_EE", !65, i64 0}
!65 = !{!"_ZTSNSt12_Vector_baseIN7hipsycl4glue4sscp63__acpp_sscp_emit_param_type_annotation_fcall_specialized_configIPKNS2_24fcall_specialized_configEEESaIS7_EE12_Vector_implE", !66, i64 0}
!66 = !{!"_ZTSNSt12_Vector_baseIN7hipsycl4glue4sscp63__acpp_sscp_emit_param_type_annotation_fcall_specialized_configIPKNS2_24fcall_specialized_configEEESaIS7_EE17_Vector_impl_dataE", !8, i64 0, !8, i64 8, !8, i64 16}
!67 = !{!"_ZTSSt6vectorISt4pairIiiESaIS1_EE", !68, i64 0}
!68 = !{!"_ZTSSt12_Vector_baseISt4pairIiiESaIS1_EE", !69, i64 0}
!69 = !{!"_ZTSNSt12_Vector_baseISt4pairIiiESaIS1_EE12_Vector_implE", !70, i64 0}
!70 = !{!"_ZTSNSt12_Vector_baseISt4pairIiiESaIS1_EE17_Vector_impl_dataE", !8, i64 0, !8, i64 8, !8, i64 16}
!71 = !{!"_ZTSSt5arrayImLm2EE", !9, i64 0}
!72 = !{!"_ZTSSt13unordered_mapINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEmSt4hashIS5_ESt8equal_toIS5_ESaISt4pairIKS5_mEEE", !73, i64 0}
!73 = !{!"_ZTSSt10_HashtableINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESt4pairIKS5_mESaIS8_ENSt8__detail10_Select1stESt8equal_toIS5_ESt4hashIS5_ENSA_18_Mod_range_hashingENSA_20_Default_ranged_hashENSA_20_Prime_rehash_policyENSA_17_Hashtable_traitsILb1ELb0ELb1EEEE", !8, i64 0, !11, i64 8, !74, i64 16, !11, i64 24, !75, i64 32, !8, i64 48}
!74 = !{!"_ZTSNSt8__detail15_Hash_node_baseE", !8, i64 0}
!75 = !{!"_ZTSNSt8__detail20_Prime_rehash_policyE", !76, i64 0, !11, i64 8}
!76 = !{!"float", !9, i64 0}
!77 = !{!6, !11, i64 16}
!78 = !{!79, !79, i64 0}
!79 = !{!"vtable pointer", !10, i64 0}
!80 = !{!81, !81, i64 0}
!81 = !{!"int", !9, i64 0}


; Effective O3 passes: EarlyCSEPass,GVNPass,GlobalOptPass,IPSCCPPass,IndVarSimplifyPass,InferFunctionAttrsPass,InstCombinePass,LoopSimplifyPass,SROAPass,SimplifyCFGPass,TailCallElimPass