; 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.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.std::shared_ptr.0" = type { %"class.std::__shared_ptr.1" }
%"class.std::__shared_ptr.1" = type { ptr, %"class.std::__shared_count" }
%"class.std::__shared_count" = type { ptr }
%"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.cl::Event" = type { %"class.cl::detail::Wrapper.61" }
%"class.cl::detail::Wrapper.61" = type { ptr }
%"class.std::allocator.76" = type { i8 }
%"class.cl::Context" = type { %"class.cl::detail::Wrapper.51" }
%"class.cl::detail::Wrapper.51" = type { ptr }
%"class.hipsycl::rt::device_id" = type { %"struct.hipsycl::rt::backend_descriptor", i32 }
%"struct.hipsycl::rt::backend_descriptor" = type { i32, i32, i32 }
%"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.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.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
@.str.4 = external hidden unnamed_addr constant [47 x i8], align 1
@__func__._ZN7hipsycl2rt9ocl_queue21submit_queue_wait_forERKSt10shared_ptrINS0_8dag_nodeEE = external hidden unnamed_addr constant [22 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

; 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 #1 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 #2 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 #2 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 #2 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 #2 align 2

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

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

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

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

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

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

; Function Attrs: mustprogress uwtable
declare dso_local void @_ZN2cl5EventC2Ev(ptr noundef nonnull align 8 dereferenceable(8)) unnamed_addr #2 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) #2 align 2

; Function Attrs: mustprogress uwtable
declare dso_local void @_ZN7hipsycl2rt9ocl_queue21register_submitted_opEN2cl5EventE(ptr noundef nonnull align 8 dereferenceable(584), ptr noundef) #2 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 #4 align 2

; Function Attrs: mustprogress nounwind uwtable
declare dso_local void @_ZNSt6vectorIN2cl5EventESaIS1_EED2Ev(ptr noundef nonnull align 8 dereferenceable(24)) unnamed_addr #1 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)) #4

; 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) #3

; Function Attrs: mustprogress uwtable
define dso_local void @_ZN7hipsycl2rt9ocl_queue21submit_queue_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 #2 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 %"class.std::shared_ptr.0", align 8
  %9 = alloca %"class.std::vector.74", align 8
  %10 = alloca %"class.std::initializer_list", align 8
  %11 = alloca [1 x %"class.cl::Event"], align 8
  %12 = alloca ptr, align 8
  %13 = alloca ptr, align 8
  %14 = alloca i32, align 4
  %15 = alloca %"class.std::allocator.76", align 1
  %16 = alloca %"class.cl::Context", align 8
  %17 = alloca %"class.hipsycl::rt::device_id", align 4
  %18 = alloca %"class.cl::Context", align 8
  %19 = alloca %"class.hipsycl::rt::device_id", align 4
  %20 = alloca i32, align 4
  %21 = alloca %"class.cl::Event", align 8
  %22 = alloca i32, align 4
  %23 = alloca %"class.hipsycl::rt::source_location", align 8
  %24 = alloca %"class.std::__cxx11::basic_string", align 8
  %25 = alloca %"class.std::allocator.58", align 1
  %26 = alloca %"class.std::__cxx11::basic_string", align 8
  %27 = alloca %"class.std::allocator.58", align 1
  %28 = alloca %"class.hipsycl::rt::error_info", align 8
  %29 = alloca %"class.std::__cxx11::basic_string", align 8
  %30 = alloca %"class.std::allocator.58", align 1
  %31 = alloca %"class.hipsycl::rt::error_code", align 8
  %32 = alloca %"class.std::__cxx11::basic_string", align 8
  %33 = alloca %"class.std::allocator.58", align 1
  %34 = alloca %"class.cl::Event", 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
  %35 = load ptr, ptr %5, align 8
  call void @llvm.lifetime.start.p0(i64 8, ptr %7) #6
  call void @llvm.lifetime.start.p0(i64 16, ptr %8) #6
  %36 = load ptr, ptr %6, align 8, !tbaa !5
  %37 = call noundef ptr @_ZNKSt19__shared_ptr_accessIN7hipsycl2rt8dag_nodeELN9__gnu_cxx12_Lock_policyE2ELb0ELb0EEptEv(ptr noundef nonnull align 1 dereferenceable(1) %36) #6
  call void @_ZNK7hipsycl2rt8dag_node9get_eventEv(ptr dead_on_unwind writable sret(%"class.std::shared_ptr.0") align 8 %8, ptr noundef nonnull align 8 dereferenceable(384) %37)
  %38 = call noundef ptr @_ZNKSt12__shared_ptrIN7hipsycl2rt14dag_node_eventELN9__gnu_cxx12_Lock_policyE2EE3getEv(ptr noundef nonnull align 8 dereferenceable(16) %8) #6
  call void @_ZNSt12__shared_ptrIN7hipsycl2rt14dag_node_eventELN9__gnu_cxx12_Lock_policyE2EED2Ev(ptr noundef nonnull align 8 dereferenceable(16) %8) #6
  call void @llvm.lifetime.end.p0(i64 16, ptr %8) #6
  store ptr %38, ptr %7, align 8, !tbaa !5
  call void @llvm.lifetime.start.p0(i64 24, ptr %9) #6
  call void @llvm.lifetime.start.p0(i64 8, ptr %11) #6
  store ptr %11, ptr %12, align 8
  %39 = load ptr, ptr %7, align 8, !tbaa !5
  invoke void @_ZNK7hipsycl2rt14ocl_node_event9get_eventEv(ptr dead_on_unwind writable sret(%"class.cl::Event") align 8 %11, ptr noundef nonnull align 8 dereferenceable(40) %39)
          to label %40 unwind label %97

40:                                               ; preds = %3
  %41 = getelementptr inbounds %"class.std::initializer_list", ptr %10, i32 0, i32 0
  store ptr %11, ptr %41, align 8, !tbaa !9
  %42 = getelementptr inbounds %"class.std::initializer_list", ptr %10, i32 0, i32 1
  store i64 1, ptr %42, align 8, !tbaa !12
  call void @llvm.lifetime.start.p0(i64 1, ptr %15) #6
  call void @_ZNSaIN2cl5EventEEC2Ev(ptr noundef nonnull align 1 dereferenceable(1) %15) #6
  %43 = getelementptr inbounds { ptr, i64 }, ptr %10, i32 0, i32 0
  %44 = load ptr, ptr %43, align 8
  %45 = getelementptr inbounds { ptr, i64 }, ptr %10, i32 0, i32 1
  %46 = load i64, ptr %45, align 8
  invoke void @_ZNSt6vectorIN2cl5EventESaIS1_EEC2ESt16initializer_listIS1_ERKS2_(ptr noundef nonnull align 8 dereferenceable(24) %9, ptr %44, i64 %46, ptr noundef nonnull align 1 dereferenceable(1) %15)
          to label %47 unwind label %108

47:                                               ; preds = %40
  call void @_ZNSt15__new_allocatorIN2cl5EventEED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %15) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr %15) #6
  %48 = getelementptr inbounds [1 x %"class.cl::Event"], ptr %11, i32 0, i32 0
  %49 = getelementptr inbounds %"class.cl::Event", ptr %48, i64 1
  br label %50

50:                                               ; preds = %50, %47
  %51 = phi ptr [ %49, %47 ], [ %52, %50 ]
  %52 = getelementptr inbounds %"class.cl::Event", ptr %51, i64 -1
  call void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %52) #6
  %53 = icmp eq ptr %52, %48
  br i1 %53, label %54, label %50

54:                                               ; preds = %50
  call void @llvm.lifetime.end.p0(i64 8, ptr %11) #6
  call void @llvm.lifetime.start.p0(i64 8, ptr %16) #6
  %55 = getelementptr inbounds %"class.hipsycl::rt::ocl_queue", ptr %35, i32 0, i32 1
  %56 = load ptr, ptr %55, align 8, !tbaa !13
  %57 = load ptr, ptr %7, align 8, !tbaa !5
  %58 = invoke { i64, i64 } @_ZNK7hipsycl2rt14ocl_node_event10get_deviceEv(ptr noundef nonnull align 8 dereferenceable(40) %57)
          to label %59 unwind label %120

59:                                               ; preds = %54
  %60 = getelementptr inbounds { i64, i64 }, ptr %17, i32 0, i32 0
  %61 = extractvalue { i64, i64 } %58, 0
  store i64 %61, ptr %60, align 4
  %62 = getelementptr inbounds { i64, i64 }, ptr %17, i32 0, i32 1
  %63 = extractvalue { i64, i64 } %58, 1
  store i64 %63, ptr %62, align 4
  %64 = getelementptr inbounds { i64, i64 }, ptr %17, i32 0, i32 0
  %65 = load i64, ptr %64, align 4
  %66 = getelementptr inbounds { i64, i64 }, ptr %17, i32 0, i32 1
  %67 = load i64, ptr %66, align 4
  invoke void @_ZN7hipsycl2rt20ocl_hardware_manager11get_contextENS0_9device_idE(ptr dead_on_unwind writable sret(%"class.cl::Context") align 8 %16, ptr noundef nonnull align 8 dereferenceable(60) %56, i64 %65, i64 %67)
          to label %68 unwind label %120

68:                                               ; preds = %59
  call void @llvm.lifetime.start.p0(i64 8, ptr %18) #6
  %69 = getelementptr inbounds %"class.hipsycl::rt::ocl_queue", ptr %35, i32 0, i32 1
  %70 = load ptr, ptr %69, align 8, !tbaa !13
  %71 = getelementptr inbounds %"class.hipsycl::rt::ocl_queue", ptr %35, i32 0, i32 1
  %72 = load ptr, ptr %71, align 8, !tbaa !13
  %73 = getelementptr inbounds %"class.hipsycl::rt::ocl_queue", ptr %35, i32 0, i32 2
  %74 = load i64, ptr %73, align 8, !tbaa !81
  %75 = load ptr, ptr %72, align 8, !tbaa !82
  %76 = getelementptr inbounds ptr, ptr %75, i64 3
  %77 = load ptr, ptr %76, align 8
  %78 = invoke { i64, i64 } %77(ptr noundef nonnull align 8 dereferenceable(60) %72, i64 noundef %74)
          to label %79 unwind label %124

79:                                               ; preds = %68
  %80 = getelementptr inbounds { i64, i64 }, ptr %19, i32 0, i32 0
  %81 = extractvalue { i64, i64 } %78, 0
  store i64 %81, ptr %80, align 4
  %82 = getelementptr inbounds { i64, i64 }, ptr %19, i32 0, i32 1
  %83 = extractvalue { i64, i64 } %78, 1
  store i64 %83, ptr %82, align 4
  %84 = getelementptr inbounds { i64, i64 }, ptr %19, i32 0, i32 0
  %85 = load i64, ptr %84, align 4
  %86 = getelementptr inbounds { i64, i64 }, ptr %19, i32 0, i32 1
  %87 = load i64, ptr %86, align 4
  invoke void @_ZN7hipsycl2rt20ocl_hardware_manager11get_contextENS0_9device_idE(ptr dead_on_unwind writable sret(%"class.cl::Context") align 8 %18, ptr noundef nonnull align 8 dereferenceable(60) %70, i64 %85, i64 %87)
          to label %88 unwind label %124

88:                                               ; preds = %79
  %89 = invoke noundef zeroext i1 @_ZN2cl6detailneIP11_cl_contextEEbRKNS0_7WrapperIT_EES8_(ptr noundef nonnull align 8 dereferenceable(8) %16, ptr noundef nonnull align 8 dereferenceable(8) %18)
          to label %90 unwind label %128

90:                                               ; preds = %88
  call void @_ZN2cl6detail7WrapperIP11_cl_contextED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %18) #6
  call void @llvm.lifetime.end.p0(i64 8, ptr %18) #6
  call void @_ZN2cl6detail7WrapperIP11_cl_contextED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %16) #6
  call void @llvm.lifetime.end.p0(i64 8, ptr %16) #6
  br i1 %89, label %91, label %138

91:                                               ; preds = %90
  %92 = load ptr, ptr %6, align 8, !tbaa !5
  %93 = load ptr, ptr %35, align 8, !tbaa !82
  %94 = getelementptr inbounds ptr, ptr %93, i64 7
  %95 = load ptr, ptr %94, align 8
  invoke void %95(ptr dead_on_unwind writable sret(%"class.hipsycl::rt::result") align 8 %0, ptr noundef nonnull align 8 dereferenceable(584) %35, ptr noundef nonnull align 8 dereferenceable(16) %92)
          to label %96 unwind label %134

96:                                               ; preds = %91
  store i32 1, ptr %20, align 4
  br label %213

97:                                               ; preds = %3
  %98 = landingpad { ptr, i32 }
          cleanup
  %99 = extractvalue { ptr, i32 } %98, 0
  store ptr %99, ptr %13, align 8
  %100 = extractvalue { ptr, i32 } %98, 1
  store i32 %100, ptr %14, align 4
  %101 = load ptr, ptr %12, align 8
  %102 = icmp eq ptr %11, %101
  br i1 %102, label %107, label %103

103:                                              ; preds = %103, %97
  %104 = phi ptr [ %101, %97 ], [ %105, %103 ]
  %105 = getelementptr inbounds %"class.cl::Event", ptr %104, i64 -1
  call void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %105) #6
  %106 = icmp eq ptr %105, %11
  br i1 %106, label %107, label %103

107:                                              ; preds = %103, %97
  br label %119

108:                                              ; preds = %40
  %109 = landingpad { ptr, i32 }
          cleanup
  %110 = extractvalue { ptr, i32 } %109, 0
  store ptr %110, ptr %13, align 8
  %111 = extractvalue { ptr, i32 } %109, 1
  store i32 %111, ptr %14, align 4
  call void @_ZNSt15__new_allocatorIN2cl5EventEED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %15) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr %15) #6
  %112 = getelementptr inbounds [1 x %"class.cl::Event"], ptr %11, i32 0, i32 0
  %113 = getelementptr inbounds %"class.cl::Event", ptr %112, i64 1
  br label %114

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

118:                                              ; preds = %114
  br label %119

119:                                              ; preds = %118, %107
  call void @llvm.lifetime.end.p0(i64 8, ptr %11) #6
  br label %215

120:                                              ; preds = %59, %54
  %121 = landingpad { ptr, i32 }
          cleanup
  %122 = extractvalue { ptr, i32 } %121, 0
  store ptr %122, ptr %13, align 8
  %123 = extractvalue { ptr, i32 } %121, 1
  store i32 %123, ptr %14, align 4
  br label %133

124:                                              ; preds = %79, %68
  %125 = landingpad { ptr, i32 }
          cleanup
  %126 = extractvalue { ptr, i32 } %125, 0
  store ptr %126, ptr %13, align 8
  %127 = extractvalue { ptr, i32 } %125, 1
  store i32 %127, ptr %14, align 4
  br label %132

128:                                              ; preds = %88
  %129 = landingpad { ptr, i32 }
          cleanup
  %130 = extractvalue { ptr, i32 } %129, 0
  store ptr %130, ptr %13, align 8
  %131 = extractvalue { ptr, i32 } %129, 1
  store i32 %131, ptr %14, align 4
  call void @_ZN2cl6detail7WrapperIP11_cl_contextED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %18) #6
  br label %132

132:                                              ; preds = %128, %124
  call void @llvm.lifetime.end.p0(i64 8, ptr %18) #6
  call void @_ZN2cl6detail7WrapperIP11_cl_contextED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %16) #6
  br label %133

133:                                              ; preds = %132, %120
  call void @llvm.lifetime.end.p0(i64 8, ptr %16) #6
  br label %214

134:                                              ; preds = %91
  %135 = landingpad { ptr, i32 }
          cleanup
  %136 = extractvalue { ptr, i32 } %135, 0
  store ptr %136, ptr %13, align 8
  %137 = extractvalue { ptr, i32 } %135, 1
  store i32 %137, ptr %14, align 4
  br label %214

138:                                              ; preds = %90
  call void @llvm.lifetime.start.p0(i64 8, ptr %21) #6
  invoke void @_ZN2cl5EventC2Ev(ptr noundef nonnull align 8 dereferenceable(8) %21)
          to label %139 unwind label %155

139:                                              ; preds = %138
  call void @llvm.lifetime.start.p0(i64 4, ptr %22) #6
  %140 = getelementptr inbounds %"class.hipsycl::rt::ocl_queue", ptr %35, i32 0, i32 3
  %141 = invoke noundef i32 @_ZNK2cl12CommandQueue26enqueueBarrierWithWaitListEPKSt6vectorINS_5EventESaIS2_EEPS2_(ptr noundef nonnull align 8 dereferenceable(8) %140, ptr noundef %9, ptr noundef %21)
          to label %142 unwind label %159

142:                                              ; preds = %139
  store i32 %141, ptr %22, align 4, !tbaa !84
  %143 = load i32, ptr %22, align 4, !tbaa !84
  %144 = icmp ne i32 %143, 0
  br i1 %144, label %145, label %202

145:                                              ; preds = %142
  call void @llvm.lifetime.start.p0(i64 72, ptr %23) #6
  call void @llvm.lifetime.start.p0(i64 32, ptr %24) #6
  call void @llvm.lifetime.start.p0(i64 1, ptr %25) #6
  call void @_ZNSaIcEC2Ev(ptr noundef nonnull align 1 dereferenceable(1) %25) #6
  invoke void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_(ptr noundef nonnull align 8 dereferenceable(32) %24, ptr noundef @__func__._ZN7hipsycl2rt9ocl_queue21submit_queue_wait_forERKSt10shared_ptrINS0_8dag_nodeEE, ptr noundef nonnull align 1 dereferenceable(1) %25)
          to label %146 unwind label %163

146:                                              ; preds = %145
  call void @llvm.lifetime.start.p0(i64 32, ptr %26) #6
  call void @llvm.lifetime.start.p0(i64 1, ptr %27) #6
  call void @_ZNSaIcEC2Ev(ptr noundef nonnull align 1 dereferenceable(1) %27) #6
  invoke void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_(ptr noundef nonnull align 8 dereferenceable(32) %26, ptr noundef @.str, ptr noundef nonnull align 1 dereferenceable(1) %27)
          to label %147 unwind label %167

147:                                              ; preds = %146
  invoke void @_ZN7hipsycl2rt15source_locationC2ERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEES9_i(ptr noundef nonnull align 8 dereferenceable(68) %23, ptr noundef nonnull align 8 dereferenceable(32) %24, ptr noundef nonnull align 8 dereferenceable(32) %26, i32 noundef 359)
          to label %148 unwind label %171

148:                                              ; preds = %147
  call void @llvm.lifetime.start.p0(i64 80, ptr %28) #6
  call void @llvm.lifetime.start.p0(i64 32, ptr %29) #6
  call void @llvm.lifetime.start.p0(i64 1, ptr %30) #6
  call void @_ZNSaIcEC2Ev(ptr noundef nonnull align 1 dereferenceable(1) %30) #6
  invoke void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_(ptr noundef nonnull align 8 dereferenceable(32) %29, ptr noundef @.str.4, ptr noundef nonnull align 1 dereferenceable(1) %30)
          to label %149 unwind label %175

149:                                              ; preds = %148
  call void @llvm.lifetime.start.p0(i64 32, ptr %32) #6
  call void @llvm.lifetime.start.p0(i64 1, ptr %33) #6
  call void @_ZNSaIcEC2Ev(ptr noundef nonnull align 1 dereferenceable(1) %33) #6
  invoke void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_(ptr noundef nonnull align 8 dereferenceable(32) %32, ptr noundef @.str.3, ptr noundef nonnull align 1 dereferenceable(1) %33)
          to label %150 unwind label %179

150:                                              ; preds = %149
  %151 = load i32, ptr %22, align 4, !tbaa !84
  invoke void @_ZN7hipsycl2rt10error_codeC2ERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEi(ptr noundef nonnull align 8 dereferenceable(40) %31, ptr noundef nonnull align 8 dereferenceable(32) %32, i32 noundef %151)
          to label %152 unwind label %183

152:                                              ; preds = %150
  invoke void @_ZN7hipsycl2rt10error_infoC2ERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEENS0_10error_codeENS0_10error_typeE(ptr noundef nonnull align 8 dereferenceable(76) %28, ptr noundef nonnull align 8 dereferenceable(32) %29, ptr noundef %31, i32 noundef 1)
          to label %153 unwind label %187

153:                                              ; preds = %152
  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) %23, ptr noundef nonnull align 8 dereferenceable(76) %28)
          to label %154 unwind label %191

154:                                              ; preds = %153
  call void @_ZN7hipsycl2rt10error_infoD2Ev(ptr noundef nonnull align 8 dereferenceable(76) %28) #6
  call void @_ZN7hipsycl2rt10error_codeD2Ev(ptr noundef nonnull align 8 dereferenceable(40) %31) #6
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %32) #6
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %33) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr %33) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr %32) #6
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %29) #6
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %30) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr %30) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr %29) #6
  call void @llvm.lifetime.end.p0(i64 80, ptr %28) #6
  call void @_ZN7hipsycl2rt15source_locationD2Ev(ptr noundef nonnull align 8 dereferenceable(68) %23) #6
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %26) #6
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %27) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr %27) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr %26) #6
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %24) #6
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %25) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr %25) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr %24) #6
  call void @llvm.lifetime.end.p0(i64 72, ptr %23) #6
  store i32 1, ptr %20, align 4
  br label %210

155:                                              ; preds = %138
  %156 = landingpad { ptr, i32 }
          cleanup
  %157 = extractvalue { ptr, i32 } %156, 0
  store ptr %157, ptr %13, align 8
  %158 = extractvalue { ptr, i32 } %156, 1
  store i32 %158, ptr %14, align 4
  br label %212

159:                                              ; preds = %204, %202, %139
  %160 = landingpad { ptr, i32 }
          cleanup
  %161 = extractvalue { ptr, i32 } %160, 0
  store ptr %161, ptr %13, align 8
  %162 = extractvalue { ptr, i32 } %160, 1
  store i32 %162, ptr %14, align 4
  br label %211

163:                                              ; preds = %145
  %164 = landingpad { ptr, i32 }
          cleanup
  %165 = extractvalue { ptr, i32 } %164, 0
  store ptr %165, ptr %13, align 8
  %166 = extractvalue { ptr, i32 } %164, 1
  store i32 %166, ptr %14, align 4
  br label %201

167:                                              ; preds = %146
  %168 = landingpad { ptr, i32 }
          cleanup
  %169 = extractvalue { ptr, i32 } %168, 0
  store ptr %169, ptr %13, align 8
  %170 = extractvalue { ptr, i32 } %168, 1
  store i32 %170, ptr %14, align 4
  br label %200

171:                                              ; preds = %147
  %172 = landingpad { ptr, i32 }
          cleanup
  %173 = extractvalue { ptr, i32 } %172, 0
  store ptr %173, ptr %13, align 8
  %174 = extractvalue { ptr, i32 } %172, 1
  store i32 %174, ptr %14, align 4
  br label %199

175:                                              ; preds = %148
  %176 = landingpad { ptr, i32 }
          cleanup
  %177 = extractvalue { ptr, i32 } %176, 0
  store ptr %177, ptr %13, align 8
  %178 = extractvalue { ptr, i32 } %176, 1
  store i32 %178, ptr %14, align 4
  br label %198

179:                                              ; preds = %149
  %180 = landingpad { ptr, i32 }
          cleanup
  %181 = extractvalue { ptr, i32 } %180, 0
  store ptr %181, ptr %13, align 8
  %182 = extractvalue { ptr, i32 } %180, 1
  store i32 %182, ptr %14, align 4
  br label %197

183:                                              ; preds = %150
  %184 = landingpad { ptr, i32 }
          cleanup
  %185 = extractvalue { ptr, i32 } %184, 0
  store ptr %185, ptr %13, align 8
  %186 = extractvalue { ptr, i32 } %184, 1
  store i32 %186, ptr %14, align 4
  br label %196

187:                                              ; preds = %152
  %188 = landingpad { ptr, i32 }
          cleanup
  %189 = extractvalue { ptr, i32 } %188, 0
  store ptr %189, ptr %13, align 8
  %190 = extractvalue { ptr, i32 } %188, 1
  store i32 %190, ptr %14, align 4
  br label %195

191:                                              ; preds = %153
  %192 = landingpad { ptr, i32 }
          cleanup
  %193 = extractvalue { ptr, i32 } %192, 0
  store ptr %193, ptr %13, align 8
  %194 = extractvalue { ptr, i32 } %192, 1
  store i32 %194, ptr %14, align 4
  call void @_ZN7hipsycl2rt10error_infoD2Ev(ptr noundef nonnull align 8 dereferenceable(76) %28) #6
  br label %195

195:                                              ; preds = %191, %187
  call void @_ZN7hipsycl2rt10error_codeD2Ev(ptr noundef nonnull align 8 dereferenceable(40) %31) #6
  br label %196

196:                                              ; preds = %195, %183
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %32) #6
  br label %197

197:                                              ; preds = %196, %179
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %33) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr %33) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr %32) #6
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %29) #6
  br label %198

198:                                              ; preds = %197, %175
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %30) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr %30) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr %29) #6
  call void @llvm.lifetime.end.p0(i64 80, ptr %28) #6
  call void @_ZN7hipsycl2rt15source_locationD2Ev(ptr noundef nonnull align 8 dereferenceable(68) %23) #6
  br label %199

199:                                              ; preds = %198, %171
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %26) #6
  br label %200

200:                                              ; preds = %199, %167
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %27) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr %27) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr %26) #6
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %24) #6
  br label %201

201:                                              ; preds = %200, %163
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %25) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr %25) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr %24) #6
  call void @llvm.lifetime.end.p0(i64 72, ptr %23) #6
  br label %211

202:                                              ; preds = %142
  invoke void @_ZN2cl5EventC2ERKS0_(ptr noundef nonnull align 8 dereferenceable(8) %34, ptr noundef nonnull align 8 dereferenceable(8) %21)
          to label %203 unwind label %159

203:                                              ; preds = %202
  invoke void @_ZN7hipsycl2rt9ocl_queue21register_submitted_opEN2cl5EventE(ptr noundef nonnull align 8 dereferenceable(584) %35, ptr noundef %34)
          to label %204 unwind label %206

204:                                              ; preds = %203
  call void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %34) #6
  invoke void @_ZN7hipsycl2rt12make_successEv(ptr dead_on_unwind writable sret(%"class.hipsycl::rt::result") align 8 %0)
          to label %205 unwind label %159

205:                                              ; preds = %204
  store i32 1, ptr %20, align 4
  br label %210

206:                                              ; preds = %203
  %207 = landingpad { ptr, i32 }
          cleanup
  %208 = extractvalue { ptr, i32 } %207, 0
  store ptr %208, ptr %13, align 8
  %209 = extractvalue { ptr, i32 } %207, 1
  store i32 %209, ptr %14, align 4
  call void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %34) #6
  br label %211

210:                                              ; preds = %205, %154
  call void @llvm.lifetime.end.p0(i64 4, ptr %22) #6
  call void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %21) #6
  call void @llvm.lifetime.end.p0(i64 8, ptr %21) #6
  br label %213

211:                                              ; preds = %206, %201, %159
  call void @llvm.lifetime.end.p0(i64 4, ptr %22) #6
  call void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %21) #6
  br label %212

212:                                              ; preds = %211, %155
  call void @llvm.lifetime.end.p0(i64 8, ptr %21) #6
  br label %214

213:                                              ; preds = %210, %96
  call void @_ZNSt6vectorIN2cl5EventESaIS1_EED2Ev(ptr noundef nonnull align 8 dereferenceable(24) %9) #6
  call void @llvm.lifetime.end.p0(i64 24, ptr %9) #6
  call void @llvm.lifetime.end.p0(i64 8, ptr %7) #6
  ret void

214:                                              ; preds = %212, %134, %133
  call void @_ZNSt6vectorIN2cl5EventESaIS1_EED2Ev(ptr noundef nonnull align 8 dereferenceable(24) %9) #6
  br label %215

215:                                              ; preds = %214, %119
  call void @llvm.lifetime.end.p0(i64 24, ptr %9) #6
  call void @llvm.lifetime.end.p0(i64 8, ptr %7) #6
  br label %216

216:                                              ; preds = %215
  %217 = load ptr, ptr %13, align 8
  %218 = load i32, ptr %14, align 4
  %219 = insertvalue { ptr, i32 } poison, ptr %217, 0
  %220 = insertvalue { ptr, i32 } %219, i32 %218, 1
  resume { ptr, i32 } %220
}

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

declare void @_ZNK7hipsycl2rt8dag_node9get_eventEv(ptr dead_on_unwind writable sret(%"class.std::shared_ptr.0") align 8, ptr noundef nonnull align 8 dereferenceable(384)) #5

; Function Attrs: mustprogress nounwind uwtable
declare dso_local noundef ptr @_ZNKSt12__shared_ptrIN7hipsycl2rt14dag_node_eventELN9__gnu_cxx12_Lock_policyE2EE3getEv(ptr noundef nonnull align 8 dereferenceable(16)) #1 align 2

declare void @_ZNK7hipsycl2rt14ocl_node_event9get_eventEv(ptr dead_on_unwind writable sret(%"class.cl::Event") align 8, ptr noundef nonnull align 8 dereferenceable(40)) #5

; Function Attrs: mustprogress nounwind uwtable
declare dso_local void @_ZNSaIN2cl5EventEEC2Ev(ptr noundef nonnull align 1 dereferenceable(1)) unnamed_addr #1 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 #2 align 2

; Function Attrs: inlinehint mustprogress uwtable
declare dso_local noundef zeroext i1 @_ZN2cl6detailneIP11_cl_contextEEbRKNS0_7WrapperIT_EES8_(ptr noundef nonnull align 8 dereferenceable(8), ptr noundef nonnull align 8 dereferenceable(8)) #4

declare void @_ZN7hipsycl2rt20ocl_hardware_manager11get_contextENS0_9device_idE(ptr dead_on_unwind writable sret(%"class.cl::Context") align 8, ptr noundef nonnull align 8 dereferenceable(60), i64, i64) #5

declare { i64, i64 } @_ZNK7hipsycl2rt14ocl_node_event10get_deviceEv(ptr noundef nonnull align 8 dereferenceable(40)) #5

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

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

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

attributes #0 = { nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) }
attributes #1 = { 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 #2 = { 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 #3 = { 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 #4 = { 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 #5 = { "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 0}
!10 = !{!"_ZTSSt16initializer_listIN2cl5EventEE", !6, i64 0, !11, i64 8}
!11 = !{!"long", !7, i64 0}
!12 = !{!10, !11, i64 8}
!13 = !{!14, !6, i64 8}
!14 = !{!"_ZTSN7hipsycl2rt9ocl_queueE", !15, i64 0, !6, i64 8, !11, i64 16, !16, i64 24, !18, i64 32, !20, i64 48, !36, i64 232, !39, i64 248, !42, i64 304, !45, i64 312, !54, i64 368, !76, i64 528}
!15 = !{!"_ZTSN7hipsycl2rt13inorder_queueE"}
!16 = !{!"_ZTSN2cl12CommandQueueE", !17, i64 0}
!17 = !{!"_ZTSN2cl6detail7WrapperIP17_cl_command_queueEE", !6, i64 0}
!18 = !{!"_ZTSN7hipsycl2rt28ocl_sscp_code_object_invokerE", !19, i64 0, !6, i64 8}
!19 = !{!"_ZTSN7hipsycl2rt24sscp_code_object_invokerE"}
!20 = !{!"_ZTSN7hipsycl2rt13worker_threadE", !21, i64 0, !23, i64 8, !26, i64 16, !28, i64 64, !30, i64 104}
!21 = !{!"_ZTSSt6thread", !22, i64 0}
!22 = !{!"_ZTSNSt6thread2idE", !11, i64 0}
!23 = !{!"_ZTSSt6atomicIbE", !24, i64 0}
!24 = !{!"_ZTSSt13__atomic_baseIbE", !25, i64 0}
!25 = !{!"bool", !7, i64 0}
!26 = !{!"_ZTSSt18condition_variable", !27, i64 0}
!27 = !{!"_ZTSSt9__condvar", !7, i64 0}
!28 = !{!"_ZTSSt5mutex", !29, i64 0}
!29 = !{!"_ZTSSt12__mutex_base", !7, i64 0}
!30 = !{!"_ZTSSt5queueISt8functionIFvvEESt5dequeIS2_SaIS2_EEE", !31, i64 0}
!31 = !{!"_ZTSSt5dequeISt8functionIFvvEESaIS2_EE", !32, i64 0}
!32 = !{!"_ZTSSt11_Deque_baseISt8functionIFvvEESaIS2_EE", !33, i64 0}
!33 = !{!"_ZTSNSt11_Deque_baseISt8functionIFvvEESaIS2_EE11_Deque_implE", !34, i64 0}
!34 = !{!"_ZTSNSt11_Deque_baseISt8functionIFvvEESaIS2_EE16_Deque_impl_dataE", !6, i64 0, !11, i64 8, !35, i64 16, !35, i64 48}
!35 = !{!"_ZTSSt15_Deque_iteratorISt8functionIFvvEERS2_PS2_E", !6, i64 0, !6, i64 8, !6, i64 16, !6, i64 24}
!36 = !{!"_ZTSSt10shared_ptrIN7hipsycl2rt12kernel_cacheEE", !37, i64 0}
!37 = !{!"_ZTSSt12__shared_ptrIN7hipsycl2rt12kernel_cacheELN9__gnu_cxx12_Lock_policyE2EE", !6, i64 0, !38, i64 8}
!38 = !{!"_ZTSSt14__shared_countILN9__gnu_cxx12_Lock_policyE2EE", !6, i64 0}
!39 = !{!"_ZTSN7hipsycl2rt9ocl_queue15protected_stateE", !40, i64 0, !28, i64 16}
!40 = !{!"_ZTSSt10shared_ptrIN7hipsycl2rt14dag_node_eventEE", !41, i64 0}
!41 = !{!"_ZTSSt12__shared_ptrIN7hipsycl2rt14dag_node_eventELN9__gnu_cxx12_Lock_policyE2EE", !6, i64 0, !38, i64 8}
!42 = !{!"_ZTSN7hipsycl6common9spin_lockE", !43, i64 0}
!43 = !{!"_ZTSSt11atomic_flag", !44, i64 0}
!44 = !{!"_ZTSSt18__atomic_flag_base", !25, i64 0}
!45 = !{!"_ZTSN7hipsycl4glue3jit19cxx_argument_mapperE", !25, i64 0, !46, i64 8, !50, i64 32}
!46 = !{!"_ZTSSt6vectorIPvSaIS0_EE", !47, i64 0}
!47 = !{!"_ZTSSt12_Vector_baseIPvSaIS0_EE", !48, i64 0}
!48 = !{!"_ZTSNSt12_Vector_baseIPvSaIS0_EE12_Vector_implE", !49, i64 0}
!49 = !{!"_ZTSNSt12_Vector_baseIPvSaIS0_EE17_Vector_impl_dataE", !6, i64 0, !6, i64 8, !6, i64 16}
!50 = !{!"_ZTSSt6vectorImSaImEE", !51, i64 0}
!51 = !{!"_ZTSSt12_Vector_baseImSaImEE", !52, i64 0}
!52 = !{!"_ZTSNSt12_Vector_baseImSaImEE12_Vector_implE", !53, i64 0}
!53 = !{!"_ZTSNSt12_Vector_baseImSaImEE17_Vector_impl_dataE", !6, i64 0, !6, i64 8, !6, i64 16}
!54 = !{!"_ZTSN7hipsycl2rt20kernel_configurationE", !55, i64 0, !59, i64 24, !63, i64 48, !67, i64 72, !50, i64 96, !71, i64 120, !75, i64 144}
!55 = !{!"_ZTSSt6vectorIN7hipsycl2rt17kernel_build_flagESaIS2_EE", !56, i64 0}
!56 = !{!"_ZTSSt12_Vector_baseIN7hipsycl2rt17kernel_build_flagESaIS2_EE", !57, i64 0}
!57 = !{!"_ZTSNSt12_Vector_baseIN7hipsycl2rt17kernel_build_flagESaIS2_EE12_Vector_implE", !58, i64 0}
!58 = !{!"_ZTSNSt12_Vector_baseIN7hipsycl2rt17kernel_build_flagESaIS2_EE17_Vector_impl_dataE", !6, i64 0, !6, i64 8, !6, i64 16}
!59 = !{!"_ZTSSt6vectorISt4pairIN7hipsycl2rt19kernel_build_optionENS2_20kernel_configuration13int_or_stringEESaIS6_EE", !60, i64 0}
!60 = !{!"_ZTSSt12_Vector_baseISt4pairIN7hipsycl2rt19kernel_build_optionENS2_20kernel_configuration13int_or_stringEESaIS6_EE", !61, i64 0}
!61 = !{!"_ZTSNSt12_Vector_baseISt4pairIN7hipsycl2rt19kernel_build_optionENS2_20kernel_configuration13int_or_stringEESaIS6_EE12_Vector_implE", !62, i64 0}
!62 = !{!"_ZTSNSt12_Vector_baseISt4pairIN7hipsycl2rt19kernel_build_optionENS2_20kernel_configuration13int_or_stringEESaIS6_EE17_Vector_impl_dataE", !6, i64 0, !6, i64 8, !6, i64 16}
!63 = !{!"_ZTSSt6vectorISt4pairIimESaIS1_EE", !64, i64 0}
!64 = !{!"_ZTSSt12_Vector_baseISt4pairIimESaIS1_EE", !65, i64 0}
!65 = !{!"_ZTSNSt12_Vector_baseISt4pairIimESaIS1_EE12_Vector_implE", !66, i64 0}
!66 = !{!"_ZTSNSt12_Vector_baseISt4pairIimESaIS1_EE17_Vector_impl_dataE", !6, i64 0, !6, i64 8, !6, i64 16}
!67 = !{!"_ZTSSt6vectorIN7hipsycl4glue4sscp63__acpp_sscp_emit_param_type_annotation_fcall_specialized_configIPKNS2_24fcall_specialized_configEEESaIS7_EE", !68, i64 0}
!68 = !{!"_ZTSSt12_Vector_baseIN7hipsycl4glue4sscp63__acpp_sscp_emit_param_type_annotation_fcall_specialized_configIPKNS2_24fcall_specialized_configEEESaIS7_EE", !69, i64 0}
!69 = !{!"_ZTSNSt12_Vector_baseIN7hipsycl4glue4sscp63__acpp_sscp_emit_param_type_annotation_fcall_specialized_configIPKNS2_24fcall_specialized_configEEESaIS7_EE12_Vector_implE", !70, i64 0}
!70 = !{!"_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}
!71 = !{!"_ZTSSt6vectorISt4pairIiiESaIS1_EE", !72, i64 0}
!72 = !{!"_ZTSSt12_Vector_baseISt4pairIiiESaIS1_EE", !73, i64 0}
!73 = !{!"_ZTSNSt12_Vector_baseISt4pairIiiESaIS1_EE12_Vector_implE", !74, i64 0}
!74 = !{!"_ZTSNSt12_Vector_baseISt4pairIiiESaIS1_EE17_Vector_impl_dataE", !6, i64 0, !6, i64 8, !6, i64 16}
!75 = !{!"_ZTSSt5arrayImLm2EE", !7, i64 0}
!76 = !{!"_ZTSSt13unordered_mapINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEmSt4hashIS5_ESt8equal_toIS5_ESaISt4pairIKS5_mEEE", !77, i64 0}
!77 = !{!"_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, !11, i64 8, !78, i64 16, !11, i64 24, !79, i64 32, !6, i64 48}
!78 = !{!"_ZTSNSt8__detail15_Hash_node_baseE", !6, i64 0}
!79 = !{!"_ZTSNSt8__detail20_Prime_rehash_policyE", !80, i64 0, !11, i64 8}
!80 = !{!"float", !7, i64 0}
!81 = !{!14, !11, i64 16}
!82 = !{!83, !83, i64 0}
!83 = !{!"vtable pointer", !8, i64 0}
!84 = !{!85, !85, i64 0}
!85 = !{!"int", !7, i64 0}


; 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.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.std::shared_ptr.0" = type { %"class.std::__shared_ptr.1" }
%"class.std::__shared_ptr.1" = type { ptr, %"class.std::__shared_count" }
%"class.std::__shared_count" = type { ptr }
%"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.cl::Event" = type { %"class.cl::detail::Wrapper.61" }
%"class.cl::detail::Wrapper.61" = type { ptr }
%"class.std::allocator.76" = type { i8 }
%"class.cl::Context" = type { %"class.cl::detail::Wrapper.51" }
%"class.cl::detail::Wrapper.51" = 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 }

@.str = external hidden unnamed_addr constant [87 x i8], align 1
@.str.3 = external hidden unnamed_addr constant [3 x i8], align 1
@.str.4 = external hidden unnamed_addr constant [47 x i8], align 1
@__func__._ZN7hipsycl2rt9ocl_queue21submit_queue_wait_forERKSt10shared_ptrINS0_8dag_nodeEE = external hidden unnamed_addr constant [22 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

; 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 #1 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 #2 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 #2 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 #2 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 #2 align 2

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

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

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

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

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

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

; Function Attrs: mustprogress uwtable
declare dso_local void @_ZN2cl5EventC2Ev(ptr noundef nonnull align 8 dereferenceable(8)) unnamed_addr #2 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 #2 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 #2 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 #4 align 2

; Function Attrs: mustprogress nounwind uwtable
declare dso_local void @_ZNSt6vectorIN2cl5EventESaIS1_EED2Ev(ptr noundef nonnull align 8 dereferenceable(24)) unnamed_addr #1 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 #4

; 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 #3

; Function Attrs: mustprogress uwtable
define dso_local void @_ZN7hipsycl2rt9ocl_queue21submit_queue_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 #2 align 2 personality ptr @__gxx_personality_v0 {
  %4 = alloca %"class.std::shared_ptr.0", align 8
  %5 = alloca %"class.std::vector.74", align 8
  %6 = alloca [1 x %"class.cl::Event"], align 8
  %7 = alloca %"class.std::allocator.76", align 1
  %8 = alloca %"class.cl::Context", align 8
  %9 = alloca %"class.cl::Context", align 8
  %10 = alloca %"class.cl::Event", align 8
  %11 = alloca %"class.hipsycl::rt::source_location", align 8
  %12 = alloca %"class.std::__cxx11::basic_string", align 8
  %13 = alloca %"class.std::allocator.58", align 1
  %14 = alloca %"class.std::__cxx11::basic_string", align 8
  %15 = alloca %"class.std::allocator.58", align 1
  %16 = alloca %"class.hipsycl::rt::error_info", align 8
  %17 = alloca %"class.std::__cxx11::basic_string", align 8
  %18 = alloca %"class.std::allocator.58", align 1
  %19 = alloca %"class.hipsycl::rt::error_code", align 8
  %20 = alloca %"class.std::__cxx11::basic_string", align 8
  %21 = alloca %"class.std::allocator.58", align 1
  %22 = alloca %"class.cl::Event", align 8
  call void @llvm.lifetime.start.p0(i64 16, ptr nonnull %4) #6
  %23 = tail call noundef ptr @_ZNKSt19__shared_ptr_accessIN7hipsycl2rt8dag_nodeELN9__gnu_cxx12_Lock_policyE2ELb0ELb0EEptEv(ptr noundef nonnull align 1 dereferenceable(1) %2) #6
  call void @_ZNK7hipsycl2rt8dag_node9get_eventEv(ptr dead_on_unwind nonnull writable sret(%"class.std::shared_ptr.0") align 8 %4, ptr noundef nonnull align 8 dereferenceable(384) %23)
  %24 = call noundef ptr @_ZNKSt12__shared_ptrIN7hipsycl2rt14dag_node_eventELN9__gnu_cxx12_Lock_policyE2EE3getEv(ptr noundef nonnull align 8 dereferenceable(16) %4) #6
  call void @_ZNSt12__shared_ptrIN7hipsycl2rt14dag_node_eventELN9__gnu_cxx12_Lock_policyE2EED2Ev(ptr noundef nonnull align 8 dereferenceable(16) %4) #6
  call void @llvm.lifetime.end.p0(i64 16, ptr nonnull %4) #6
  call void @llvm.lifetime.start.p0(i64 24, ptr nonnull %5) #6
  call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %6) #6
  invoke void @_ZNK7hipsycl2rt14ocl_node_event9get_eventEv(ptr dead_on_unwind nonnull writable sret(%"class.cl::Event") align 8 %6, ptr noundef nonnull align 8 dereferenceable(40) %24)
          to label %25 unwind label %51

25:                                               ; preds = %3
  call void @llvm.lifetime.start.p0(i64 1, ptr nonnull %7) #6
  call void @_ZNSaIN2cl5EventEEC2Ev(ptr noundef nonnull align 1 dereferenceable(1) %7) #6
  invoke void @_ZNSt6vectorIN2cl5EventESaIS1_EEC2ESt16initializer_listIS1_ERKS2_(ptr noundef nonnull align 8 dereferenceable(24) %5, ptr nonnull %6, i64 1, ptr noundef nonnull align 1 dereferenceable(1) %7)
          to label %26 unwind label %.loopexit.loopexit

26:                                               ; preds = %25
  call void @_ZNSt15__new_allocatorIN2cl5EventEED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %7) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr nonnull %7) #6
  call void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %6) #6
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %6) #6
  call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %8) #6
  %27 = getelementptr inbounds i8, ptr %1, i64 8
  %28 = load ptr, ptr %27, align 8, !tbaa !5
  %29 = invoke { i64, i64 } @_ZNK7hipsycl2rt14ocl_node_event10get_deviceEv(ptr noundef nonnull align 8 dereferenceable(40) %24)
          to label %30 unwind label %54

30:                                               ; preds = %26
  %31 = extractvalue { i64, i64 } %29, 0
  %32 = extractvalue { i64, i64 } %29, 1
  invoke void @_ZN7hipsycl2rt20ocl_hardware_manager11get_contextENS0_9device_idE(ptr dead_on_unwind nonnull writable sret(%"class.cl::Context") align 8 %8, ptr noundef nonnull align 8 dereferenceable(60) %28, i64 %31, i64 %32)
          to label %33 unwind label %54

33:                                               ; preds = %30
  call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %9) #6
  %34 = load ptr, ptr %27, align 8, !tbaa !5
  %35 = getelementptr inbounds i8, ptr %1, i64 16
  %36 = load i64, ptr %35, align 8, !tbaa !77
  %37 = load ptr, ptr %34, align 8, !tbaa !78
  %38 = getelementptr inbounds i8, ptr %37, i64 24
  %39 = load ptr, ptr %38, align 8
  %40 = invoke { i64, i64 } %39(ptr noundef nonnull align 8 dereferenceable(60) %34, i64 noundef %36)
          to label %41 unwind label %56

41:                                               ; preds = %33
  %42 = extractvalue { i64, i64 } %40, 0
  %43 = extractvalue { i64, i64 } %40, 1
  invoke void @_ZN7hipsycl2rt20ocl_hardware_manager11get_contextENS0_9device_idE(ptr dead_on_unwind nonnull writable sret(%"class.cl::Context") align 8 %9, ptr noundef nonnull align 8 dereferenceable(60) %34, i64 %42, i64 %43)
          to label %44 unwind label %56

44:                                               ; preds = %41
  %45 = invoke noundef zeroext i1 @_ZN2cl6detailneIP11_cl_contextEEbRKNS0_7WrapperIT_EES8_(ptr noundef nonnull align 8 dereferenceable(8) %8, ptr noundef nonnull align 8 dereferenceable(8) %9)
          to label %46 unwind label %58

46:                                               ; preds = %44
  call void @_ZN2cl6detail7WrapperIP11_cl_contextED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %9) #6
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %9) #6
  call void @_ZN2cl6detail7WrapperIP11_cl_contextED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %8) #6
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %8) #6
  br i1 %45, label %47, label %64

47:                                               ; preds = %46
  %48 = load ptr, ptr %1, align 8, !tbaa !78
  %49 = getelementptr inbounds i8, ptr %48, i64 56
  %50 = load ptr, ptr %49, align 8
  invoke void %50(ptr dead_on_unwind 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)
          to label %113 unwind label %62

51:                                               ; preds = %3
  %52 = landingpad { ptr, i32 }
          cleanup
  br label %.loopexit

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

.loopexit:                                        ; preds = %.loopexit.loopexit, %51
  %.pn = phi { ptr, i32 } [ %52, %51 ], [ %53, %.loopexit.loopexit ]
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %6) #6
  br label %115

54:                                               ; preds = %30, %26
  %55 = landingpad { ptr, i32 }
          cleanup
  br label %61

56:                                               ; preds = %41, %33
  %57 = landingpad { ptr, i32 }
          cleanup
  br label %60

58:                                               ; preds = %44
  %59 = landingpad { ptr, i32 }
          cleanup
  call void @_ZN2cl6detail7WrapperIP11_cl_contextED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %9) #6
  br label %60

60:                                               ; preds = %58, %56
  %.pn30 = phi { ptr, i32 } [ %59, %58 ], [ %57, %56 ]
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %9) #6
  call void @_ZN2cl6detail7WrapperIP11_cl_contextED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %8) #6
  br label %61

61:                                               ; preds = %60, %54
  %.pn30.pn = phi { ptr, i32 } [ %.pn30, %60 ], [ %55, %54 ]
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %8) #6
  br label %114

62:                                               ; preds = %47
  %63 = landingpad { ptr, i32 }
          cleanup
  br label %114

64:                                               ; preds = %46
  call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %10) #6
  invoke void @_ZN2cl5EventC2Ev(ptr noundef nonnull align 8 dereferenceable(8) %10)
          to label %65 unwind label %78

65:                                               ; preds = %64
  %66 = getelementptr inbounds i8, ptr %1, i64 24
  %67 = invoke noundef i32 @_ZNK2cl12CommandQueue26enqueueBarrierWithWaitListEPKSt6vectorINS_5EventESaIS2_EEPS2_(ptr noundef nonnull align 8 dereferenceable(8) %66, ptr noundef nonnull %5, ptr noundef nonnull %10)
          to label %68 unwind label %80

68:                                               ; preds = %65
  %.not = icmp eq i32 %67, 0
  br i1 %.not, label %105, label %69

69:                                               ; preds = %68
  call void @llvm.lifetime.start.p0(i64 72, ptr nonnull %11) #6
  call void @llvm.lifetime.start.p0(i64 32, ptr nonnull %12) #6
  call void @llvm.lifetime.start.p0(i64 1, ptr nonnull %13) #6
  call void @_ZNSaIcEC2Ev(ptr noundef nonnull align 1 dereferenceable(1) %13) #6
  invoke void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_(ptr noundef nonnull align 8 dereferenceable(32) %12, ptr noundef nonnull @__func__._ZN7hipsycl2rt9ocl_queue21submit_queue_wait_forERKSt10shared_ptrINS0_8dag_nodeEE, ptr noundef nonnull align 1 dereferenceable(1) %13)
          to label %70 unwind label %82

70:                                               ; preds = %69
  call void @llvm.lifetime.start.p0(i64 32, ptr nonnull %14) #6
  call void @llvm.lifetime.start.p0(i64 1, ptr nonnull %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 nonnull @.str, ptr noundef nonnull align 1 dereferenceable(1) %15)
          to label %71 unwind label %84

71:                                               ; preds = %70
  invoke void @_ZN7hipsycl2rt15source_locationC2ERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEES9_i(ptr noundef nonnull align 8 dereferenceable(68) %11, ptr noundef nonnull align 8 dereferenceable(32) %12, ptr noundef nonnull align 8 dereferenceable(32) %14, i32 noundef 359)
          to label %72 unwind label %86

72:                                               ; preds = %71
  call void @llvm.lifetime.start.p0(i64 80, ptr nonnull %16) #6
  call void @llvm.lifetime.start.p0(i64 32, ptr nonnull %17) #6
  call void @llvm.lifetime.start.p0(i64 1, ptr nonnull %18) #6
  call void @_ZNSaIcEC2Ev(ptr noundef nonnull align 1 dereferenceable(1) %18) #6
  invoke void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_(ptr noundef nonnull align 8 dereferenceable(32) %17, ptr noundef nonnull @.str.4, ptr noundef nonnull align 1 dereferenceable(1) %18)
          to label %73 unwind label %88

73:                                               ; preds = %72
  call void @llvm.lifetime.start.p0(i64 32, ptr nonnull %20) #6
  call void @llvm.lifetime.start.p0(i64 1, ptr nonnull %21) #6
  call void @_ZNSaIcEC2Ev(ptr noundef nonnull align 1 dereferenceable(1) %21) #6
  invoke void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_(ptr noundef nonnull align 8 dereferenceable(32) %20, ptr noundef nonnull @.str.3, ptr noundef nonnull align 1 dereferenceable(1) %21)
          to label %74 unwind label %90

74:                                               ; preds = %73
  invoke void @_ZN7hipsycl2rt10error_codeC2ERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEi(ptr noundef nonnull align 8 dereferenceable(40) %19, ptr noundef nonnull align 8 dereferenceable(32) %20, i32 noundef %67)
          to label %75 unwind label %92

75:                                               ; preds = %74
  invoke void @_ZN7hipsycl2rt10error_infoC2ERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEENS0_10error_codeENS0_10error_typeE(ptr noundef nonnull align 8 dereferenceable(76) %16, ptr noundef nonnull align 8 dereferenceable(32) %17, ptr noundef nonnull %19, i32 noundef 1)
          to label %76 unwind label %94

76:                                               ; preds = %75
  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) %11, ptr noundef nonnull align 8 dereferenceable(76) %16)
          to label %77 unwind label %96

77:                                               ; preds = %76
  call void @_ZN7hipsycl2rt10error_infoD2Ev(ptr noundef nonnull align 8 dereferenceable(76) %16) #6
  call void @_ZN7hipsycl2rt10error_codeD2Ev(ptr noundef nonnull align 8 dereferenceable(40) %19) #6
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %20) #6
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %21) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr nonnull %21) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %20) #6
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %17) #6
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %18) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr nonnull %18) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %17) #6
  call void @llvm.lifetime.end.p0(i64 80, ptr nonnull %16) #6
  call void @_ZN7hipsycl2rt15source_locationD2Ev(ptr noundef nonnull align 8 dereferenceable(68) %11) #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 nonnull %15) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %14) #6
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %12) #6
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %13) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr nonnull %13) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %12) #6
  call void @llvm.lifetime.end.p0(i64 72, ptr nonnull %11) #6
  br label %110

78:                                               ; preds = %64
  %79 = landingpad { ptr, i32 }
          cleanup
  br label %112

80:                                               ; preds = %105, %65
  %81 = landingpad { ptr, i32 }
          cleanup
  br label %111

82:                                               ; preds = %69
  %83 = landingpad { ptr, i32 }
          cleanup
  br label %104

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

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

88:                                               ; preds = %72
  %89 = landingpad { ptr, i32 }
          cleanup
  br label %101

90:                                               ; preds = %73
  %91 = landingpad { ptr, i32 }
          cleanup
  br label %100

92:                                               ; preds = %74
  %93 = landingpad { ptr, i32 }
          cleanup
  br label %99

94:                                               ; preds = %75
  %95 = landingpad { ptr, i32 }
          cleanup
  br label %98

96:                                               ; preds = %76
  %97 = landingpad { ptr, i32 }
          cleanup
  call void @_ZN7hipsycl2rt10error_infoD2Ev(ptr noundef nonnull align 8 dereferenceable(76) %16) #6
  br label %98

98:                                               ; preds = %96, %94
  %.pn33 = phi { ptr, i32 } [ %97, %96 ], [ %95, %94 ]
  call void @_ZN7hipsycl2rt10error_codeD2Ev(ptr noundef nonnull align 8 dereferenceable(40) %19) #6
  br label %99

99:                                               ; preds = %98, %92
  %.pn33.pn = phi { ptr, i32 } [ %.pn33, %98 ], [ %93, %92 ]
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %20) #6
  br label %100

100:                                              ; preds = %99, %90
  %.pn33.pn.pn = phi { ptr, i32 } [ %.pn33.pn, %99 ], [ %91, %90 ]
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %21) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr nonnull %21) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %20) #6
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %17) #6
  br label %101

101:                                              ; preds = %100, %88
  %.pn33.pn.pn.pn = phi { ptr, i32 } [ %.pn33.pn.pn, %100 ], [ %89, %88 ]
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %18) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr nonnull %18) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %17) #6
  call void @llvm.lifetime.end.p0(i64 80, ptr nonnull %16) #6
  call void @_ZN7hipsycl2rt15source_locationD2Ev(ptr noundef nonnull align 8 dereferenceable(68) %11) #6
  br label %102

102:                                              ; preds = %101, %86
  %.pn33.pn.pn.pn.pn = phi { ptr, i32 } [ %.pn33.pn.pn.pn, %101 ], [ %87, %86 ]
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %14) #6
  br label %103

103:                                              ; preds = %102, %84
  %.pn33.pn.pn.pn.pn.pn = phi { ptr, i32 } [ %.pn33.pn.pn.pn.pn, %102 ], [ %85, %84 ]
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %15) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr nonnull %15) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %14) #6
  call void @_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev(ptr noundef nonnull align 8 dereferenceable(32) %12) #6
  br label %104

104:                                              ; preds = %103, %82
  %.pn33.pn.pn.pn.pn.pn.pn = phi { ptr, i32 } [ %.pn33.pn.pn.pn.pn.pn, %103 ], [ %83, %82 ]
  call void @_ZNSt15__new_allocatorIcED2Ev(ptr noundef nonnull align 1 dereferenceable(1) %13) #6
  call void @llvm.lifetime.end.p0(i64 1, ptr nonnull %13) #6
  call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %12) #6
  call void @llvm.lifetime.end.p0(i64 72, ptr nonnull %11) #6
  br label %111

105:                                              ; preds = %68
  invoke void @_ZN2cl5EventC2ERKS0_(ptr noundef nonnull align 8 dereferenceable(8) %22, ptr noundef nonnull align 8 dereferenceable(8) %10)
          to label %106 unwind label %80

106:                                              ; preds = %105
  invoke void @_ZN7hipsycl2rt9ocl_queue21register_submitted_opEN2cl5EventE(ptr noundef nonnull align 8 dereferenceable(584) %1, ptr noundef nonnull %22)
          to label %107 unwind label %108

107:                                              ; preds = %106
  call void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %22) #6
  call void @_ZN7hipsycl2rt12make_successEv(ptr dead_on_unwind writable sret(%"class.hipsycl::rt::result") align 8 %0)
  br label %110

108:                                              ; preds = %106
  %109 = landingpad { ptr, i32 }
          cleanup
  call void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %22) #6
  br label %111

110:                                              ; preds = %107, %77
  call void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %10) #6
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %10) #6
  br label %113

111:                                              ; preds = %108, %104, %80
  %.pn33.pn.pn.pn.pn.pn.pn.pn = phi { ptr, i32 } [ %.pn33.pn.pn.pn.pn.pn.pn, %104 ], [ %109, %108 ], [ %81, %80 ]
  call void @_ZN2cl6detail7WrapperIP9_cl_eventED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %10) #6
  br label %112

112:                                              ; preds = %111, %78
  %.pn33.pn.pn.pn.pn.pn.pn.pn.pn = phi { ptr, i32 } [ %.pn33.pn.pn.pn.pn.pn.pn.pn, %111 ], [ %79, %78 ]
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %10) #6
  br label %114

113:                                              ; preds = %47, %110
  call void @_ZNSt6vectorIN2cl5EventESaIS1_EED2Ev(ptr noundef nonnull align 8 dereferenceable(24) %5) #6
  call void @llvm.lifetime.end.p0(i64 24, ptr nonnull %5) #6
  ret void

114:                                              ; preds = %112, %62, %61
  %.pn43 = phi { ptr, i32 } [ %63, %62 ], [ %.pn33.pn.pn.pn.pn.pn.pn.pn.pn, %112 ], [ %.pn30.pn, %61 ]
  call void @_ZNSt6vectorIN2cl5EventESaIS1_EED2Ev(ptr noundef nonnull align 8 dereferenceable(24) %5) #6
  br label %115

115:                                              ; preds = %114, %.loopexit
  %.pn43.pn = phi { ptr, i32 } [ %.pn43, %114 ], [ %.pn, %.loopexit ]
  call void @llvm.lifetime.end.p0(i64 24, ptr nonnull %5) #6
  resume { ptr, i32 } %.pn43.pn
}

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

declare void @_ZNK7hipsycl2rt8dag_node9get_eventEv(ptr dead_on_unwind writable sret(%"class.std::shared_ptr.0") align 8, ptr noundef nonnull align 8 dereferenceable(384)) local_unnamed_addr #5

; Function Attrs: mustprogress nounwind uwtable
declare dso_local noundef ptr @_ZNKSt12__shared_ptrIN7hipsycl2rt14dag_node_eventELN9__gnu_cxx12_Lock_policyE2EE3getEv(ptr noundef nonnull align 8 dereferenceable(16)) local_unnamed_addr #1 align 2

declare void @_ZNK7hipsycl2rt14ocl_node_event9get_eventEv(ptr dead_on_unwind writable sret(%"class.cl::Event") align 8, ptr noundef nonnull align 8 dereferenceable(40)) local_unnamed_addr #5

; Function Attrs: mustprogress nounwind uwtable
declare dso_local void @_ZNSaIN2cl5EventEEC2Ev(ptr noundef nonnull align 1 dereferenceable(1)) unnamed_addr #1 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 #2 align 2

; Function Attrs: inlinehint mustprogress uwtable
declare dso_local noundef zeroext i1 @_ZN2cl6detailneIP11_cl_contextEEbRKNS0_7WrapperIT_EES8_(ptr noundef nonnull align 8 dereferenceable(8), ptr noundef nonnull align 8 dereferenceable(8)) local_unnamed_addr #4

declare void @_ZN7hipsycl2rt20ocl_hardware_manager11get_contextENS0_9device_idE(ptr dead_on_unwind writable sret(%"class.cl::Context") align 8, ptr noundef nonnull align 8 dereferenceable(60), i64, i64) local_unnamed_addr #5

declare { i64, i64 } @_ZNK7hipsycl2rt14ocl_node_event10get_deviceEv(ptr noundef nonnull align 8 dereferenceable(40)) local_unnamed_addr #5

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

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

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

attributes #0 = { mustprogress nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) }
attributes #1 = { 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 #2 = { 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 #3 = { 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 #4 = { 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 #5 = { "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}


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