; LLVM IR from AdaptiveCpp/AdaptiveCpp :: tests/runtime/dag_builder.ll (ZN11dag_builder13default_hints11test_methodEv)

; 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::runtime_keep_alive_token" = type { %"class.std::shared_ptr" }
%"class.std::shared_ptr" = type { %"class.std::__shared_ptr" }
%"class.std::__shared_ptr" = type { ptr, %"class.std::__shared_count" }
%"class.std::__shared_count" = type { ptr }
%"class.hipsycl::rt::execution_hints" = type <{ %"class.hipsycl::rt::hints::bind_to_device", [4 x i8], %"class.hipsycl::rt::hints::bind_to_device_group", %"class.hipsycl::rt::hints::prefer_execution_lane", %"class.hipsycl::rt::hints::node_group", %"class.hipsycl::rt::hints::coarse_grained_synchronization", [7 x i8], %"class.hipsycl::rt::hints::prefer_executor", %"class.hipsycl::rt::hints::request_instrumentation_submission_timestamp", %"class.hipsycl::rt::hints::request_instrumentation_start_timestamp", %"class.hipsycl::rt::hints::request_instrumentation_finish_timestamp", %"class.hipsycl::rt::hints::instant_execution", [4 x i8] }>
%"class.hipsycl::rt::hints::bind_to_device" = type { %"class.hipsycl::rt::hints::execution_hint", %"class.hipsycl::rt::device_id" }
%"class.hipsycl::rt::hints::execution_hint" = type { i8 }
%"class.hipsycl::rt::device_id" = type { %"struct.hipsycl::rt::backend_descriptor", i32 }
%"struct.hipsycl::rt::backend_descriptor" = type { i32, i32, i32 }
%"class.hipsycl::rt::hints::bind_to_device_group" = type { %"class.hipsycl::rt::hints::execution_hint", %"class.std::vector.0" }
%"class.std::vector.0" = type { %"struct.std::_Vector_base.1" }
%"struct.std::_Vector_base.1" = type { %"struct.std::_Vector_base<hipsycl::rt::device_id, std::allocator<hipsycl::rt::device_id>>::_Vector_impl" }
%"struct.std::_Vector_base<hipsycl::rt::device_id, std::allocator<hipsycl::rt::device_id>>::_Vector_impl" = type { %"struct.std::_Vector_base<hipsycl::rt::device_id, std::allocator<hipsycl::rt::device_id>>::_Vector_impl_data" }
%"struct.std::_Vector_base<hipsycl::rt::device_id, std::allocator<hipsycl::rt::device_id>>::_Vector_impl_data" = type { ptr, ptr, ptr }
%"class.hipsycl::rt::hints::prefer_execution_lane" = type { %"class.hipsycl::rt::hints::execution_hint", i64 }
%"class.hipsycl::rt::hints::node_group" = type { %"class.hipsycl::rt::hints::execution_hint", i64 }
%"class.hipsycl::rt::hints::coarse_grained_synchronization" = type { %"class.hipsycl::rt::hints::execution_hint" }
%"class.hipsycl::rt::hints::prefer_executor" = type { %"class.hipsycl::rt::hints::execution_hint", ptr, %"class.std::shared_ptr.5" }
%"class.std::shared_ptr.5" = type { %"class.std::__shared_ptr.6" }
%"class.std::__shared_ptr.6" = type { ptr, %"class.std::__shared_count" }
%"class.hipsycl::rt::hints::request_instrumentation_submission_timestamp" = type { %"class.hipsycl::rt::hints::execution_hint" }
%"class.hipsycl::rt::hints::request_instrumentation_start_timestamp" = type { %"class.hipsycl::rt::hints::execution_hint" }
%"class.hipsycl::rt::hints::request_instrumentation_finish_timestamp" = type { %"class.hipsycl::rt::hints::execution_hint" }
%"class.hipsycl::rt::hints::instant_execution" = type { %"class.hipsycl::rt::hints::execution_hint" }
%"class.hipsycl::rt::dag_builder" = type { %"class.std::mutex", %"class.hipsycl::rt::dag", ptr }
%"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.hipsycl::rt::dag" = type { %"class.sbo::small_vector", %"class.sbo::small_vector" }
%"class.sbo::small_vector" = type { %"class.std::vector.8" }
%"class.std::vector.8" = type { %"struct.std::_Vector_base.9" }
%"struct.std::_Vector_base.9" = type { %"struct.std::_Vector_base<std::shared_ptr<hipsycl::rt::dag_node>, sbo::small_buffer_vector_allocator<std::shared_ptr<hipsycl::rt::dag_node>>>::_Vector_impl" }
%"struct.std::_Vector_base<std::shared_ptr<hipsycl::rt::dag_node>, sbo::small_buffer_vector_allocator<std::shared_ptr<hipsycl::rt::dag_node>>>::_Vector_impl" = type { %"struct.sbo::small_buffer_vector_allocator.base", %"struct.std::_Vector_base<std::shared_ptr<hipsycl::rt::dag_node>, sbo::small_buffer_vector_allocator<std::shared_ptr<hipsycl::rt::dag_node>>>::_Vector_impl_data" }
%"struct.sbo::small_buffer_vector_allocator.base" = type { [128 x i8], i8, i8 }
%"struct.std::_Vector_base<std::shared_ptr<hipsycl::rt::dag_node>, sbo::small_buffer_vector_allocator<std::shared_ptr<hipsycl::rt::dag_node>>>::_Vector_impl_data" = type { ptr, ptr, ptr }
%"class.hipsycl::rt::requirements_list" = type { %"class.sbo::small_vector", 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.11" }
%"struct.std::_Head_base.11" = type { ptr }
%"class.hipsycl::rt::kernel_launcher" = type { %"class.sbo::small_vector.12", %"struct.hipsycl::glue::kernel_launcher_data", %"class.hipsycl::rt::kernel_configuration" }
%"class.sbo::small_vector.12" = type { %"class.std::vector.13" }
%"class.std::vector.13" = type { %"struct.std::_Vector_base.14" }
%"struct.std::_Vector_base.14" = type { %"struct.std::_Vector_base<std::unique_ptr<hipsycl::rt::backend_kernel_launcher>, sbo::small_buffer_vector_allocator<std::unique_ptr<hipsycl::rt::backend_kernel_launcher>>>::_Vector_impl" }
%"struct.std::_Vector_base<std::unique_ptr<hipsycl::rt::backend_kernel_launcher>, sbo::small_buffer_vector_allocator<std::unique_ptr<hipsycl::rt::backend_kernel_launcher>>>::_Vector_impl" = type { %"struct.sbo::small_buffer_vector_allocator.base.16", %"struct.std::_Vector_base<std::unique_ptr<hipsycl::rt::backend_kernel_launcher>, sbo::small_buffer_vector_allocator<std::unique_ptr<hipsycl::rt::backend_kernel_launcher>>>::_Vector_impl_data" }
%"struct.sbo::small_buffer_vector_allocator.base.16" = type { [64 x i8], i8, i8 }
%"struct.std::_Vector_base<std::unique_ptr<hipsycl::rt::backend_kernel_launcher>, sbo::small_buffer_vector_allocator<std::unique_ptr<hipsycl::rt::backend_kernel_launcher>>>::_Vector_impl_data" = type { ptr, ptr, ptr }
%"struct.hipsycl::glue::kernel_launcher_data" = type { ptr, i32, %"class.std::vector.17", %"class.hipsycl::rt::static_array", %"class.hipsycl::rt::static_array", i32, %"class.std::function", i64, ptr, ptr, ptr }
%"class.std::vector.17" = type { %"struct.std::_Vector_base.18" }
%"struct.std::_Vector_base.18" = type { %"struct.std::_Vector_base<unsigned char, std::allocator<unsigned char>>::_Vector_impl" }
%"struct.std::_Vector_base<unsigned char, std::allocator<unsigned char>>::_Vector_impl" = type { %"struct.std::_Vector_base<unsigned char, std::allocator<unsigned char>>::_Vector_impl_data" }
%"struct.std::_Vector_base<unsigned char, std::allocator<unsigned char>>::_Vector_impl_data" = type { ptr, ptr, ptr }
%"class.hipsycl::rt::static_array" = type { %"struct.std::array" }
%"struct.std::array" = type { [3 x i64] }
%"class.std::function" = type { %"class.std::_Function_base", ptr }
%"class.std::_Function_base" = type { %"union.std::_Any_data", ptr }
%"union.std::_Any_data" = type { %"union.std::_Nocopy_types" }
%"union.std::_Nocopy_types" = type { { i64, i64 } }
%"class.hipsycl::rt::kernel_configuration" = type { %"class.std::vector.22", %"class.std::vector.27", %"class.std::vector.32", %"class.std::vector.37", %"class.std::vector.42", %"class.std::vector.47", %"struct.std::array.52" }
%"class.std::vector.22" = type { %"struct.std::_Vector_base.23" }
%"struct.std::_Vector_base.23" = 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.27" = type { %"struct.std::_Vector_base.28" }
%"struct.std::_Vector_base.28" = 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.32" = type { %"struct.std::_Vector_base.33" }
%"struct.std::_Vector_base.33" = 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.37" = type { %"struct.std::_Vector_base.38" }
%"struct.std::_Vector_base.38" = 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.42" = type { %"struct.std::_Vector_base.43" }
%"struct.std::_Vector_base.43" = 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.std::vector.47" = type { %"struct.std::_Vector_base.48" }
%"struct.std::_Vector_base.48" = 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.52" = type { [2 x i64] }
%"class.std::shared_ptr.53" = type { %"class.std::__shared_ptr.54" }
%"class.std::__shared_ptr.54" = type { ptr, %"class.std::__shared_count" }
%"class.boost::unit_test::basic_cstring" = type { ptr, ptr }
%"class.boost::test_tools::assertion_result" = type { %"class.boost::test_tools::assertion_result::readonly_property65", %"class.boost::shared_ptr" }
%"class.boost::test_tools::assertion_result::readonly_property65" = type { %"class.boost::unit_test::readonly_property" }
%"class.boost::unit_test::readonly_property" = type { %"class.boost::unit_test::class_property" }
%"class.boost::unit_test::class_property" = type { i8 }
%"class.boost::shared_ptr" = type { ptr, %"class.boost::detail::shared_count" }
%"class.boost::detail::shared_count" = type { ptr }
%"class.boost::unit_test::lazy_ostream_impl" = type { %"class.boost::unit_test::lazy_ostream.base", ptr, ptr }
%"class.boost::unit_test::lazy_ostream.base" = type <{ ptr, i8 }>

@_ZN5boost9unit_test12_GLOBAL__N_113unit_test_logE = external hidden global ptr, align 8
@.str.3 = external hidden unnamed_addr constant [87 x i8], align 1
@.str.6 = external hidden unnamed_addr constant [12 x i8], align 1
@.str.7 = external hidden unnamed_addr constant [49 x i8], align 1
@.str.8 = external hidden unnamed_addr constant [72 x i8], align 1

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

declare i32 @__gxx_personality_v0(...)

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

; Function Attrs: mustprogress uwtable
define dso_local void @_ZN11dag_builder13default_hints11test_methodEv(ptr noundef nonnull align 1 dereferenceable(1) %0) #1 align 2 personality ptr @__gxx_personality_v0 {
  %2 = alloca ptr, align 8
  %3 = alloca %"class.hipsycl::rt::runtime_keep_alive_token", align 8
  %4 = alloca %"class.hipsycl::rt::execution_hints", align 8
  %5 = alloca %"class.hipsycl::rt::device_id", align 4
  %6 = alloca %"struct.hipsycl::rt::backend_descriptor", align 4
  %7 = alloca ptr, align 8
  %8 = alloca i32, align 4
  %9 = alloca { i64, i32 }, align 4
  %10 = alloca %"class.hipsycl::rt::hints::bind_to_device", align 4
  %11 = alloca %"class.hipsycl::rt::device_id", align 4
  %12 = alloca %"class.hipsycl::rt::dag_builder", align 8
  %13 = alloca %"class.hipsycl::rt::requirements_list", align 8
  %14 = alloca %"class.std::unique_ptr", align 8
  %15 = alloca %"class.hipsycl::rt::kernel_launcher", align 8
  %16 = alloca %"struct.hipsycl::glue::kernel_launcher_data", align 8
  %17 = alloca %"class.sbo::small_vector.12", align 8
  %18 = alloca %"class.hipsycl::rt::requirements_list", align 8
  %19 = alloca %"class.std::shared_ptr.53", align 8
  %20 = alloca %"class.std::unique_ptr", align 8
  %21 = alloca ptr, align 8
  %22 = alloca %"class.boost::unit_test::basic_cstring", align 8
  %23 = alloca %"class.boost::unit_test::basic_cstring", align 8
  %24 = alloca %"class.boost::test_tools::assertion_result", align 8
  %25 = alloca %"class.boost::unit_test::lazy_ostream_impl", align 8
  %26 = alloca %"class.boost::unit_test::basic_cstring", align 8
  %27 = alloca %"class.boost::unit_test::basic_cstring", align 8
  %28 = alloca %"class.boost::unit_test::basic_cstring", align 8
  %29 = alloca %"class.boost::unit_test::basic_cstring", align 8
  %30 = alloca %"class.boost::test_tools::assertion_result", align 8
  %31 = alloca %"class.hipsycl::rt::device_id", align 4
  %32 = alloca %"class.boost::unit_test::lazy_ostream_impl", align 8
  %33 = alloca %"class.boost::unit_test::basic_cstring", align 8
  %34 = alloca %"class.boost::unit_test::basic_cstring", align 8
  store ptr %0, ptr %2, align 8, !tbaa !5
  %35 = load ptr, ptr %2, align 8
  call void @llvm.lifetime.start.p0(i64 16, ptr %3) #8
  call void @_ZN7hipsycl2rt24runtime_keep_alive_tokenC1Ev(ptr noundef nonnull align 8 dereferenceable(16) %3)
  call void @llvm.lifetime.start.p0(i64 136, ptr %4) #8
  call void @_ZN7hipsycl2rt15execution_hintsC2Ev(ptr noundef nonnull align 8 dereferenceable(132) %4) #8
  call void @llvm.lifetime.start.p0(i64 16, ptr %5) #8
  invoke void @_ZN7hipsycl2rt18backend_descriptorC2ENS0_17hardware_platformENS0_12api_platformE(ptr noundef nonnull align 4 dereferenceable(12) %6, i32 noundef 4, i32 noundef 4)
          to label %36 unwind label %122

36:                                               ; preds = %1
  call void @llvm.memcpy.p0.p0.i64(ptr align 4 %9, ptr align 4 %6, i64 12, i1 false)
  %37 = getelementptr inbounds { i64, i32 }, ptr %9, i32 0, i32 0
  %38 = load i64, ptr %37, align 4
  %39 = getelementptr inbounds { i64, i32 }, ptr %9, i32 0, i32 1
  %40 = load i32, ptr %39, align 4
  invoke void @_ZN7hipsycl2rt9device_idC1ENS0_18backend_descriptorEi(ptr noundef nonnull align 4 dereferenceable(16) %5, i64 %38, i32 %40, i32 noundef 12345)
          to label %41 unwind label %122

41:                                               ; preds = %36
  call void @llvm.lifetime.start.p0(i64 20, ptr %10) #8
  call void @llvm.memcpy.p0.p0.i64(ptr align 4 %11, ptr align 4 %5, i64 16, i1 false), !tbaa.struct !9
  %42 = getelementptr inbounds { i64, i64 }, ptr %11, i32 0, i32 0
  %43 = load i64, ptr %42, align 4
  %44 = getelementptr inbounds { i64, i64 }, ptr %11, i32 0, i32 1
  %45 = load i64, ptr %44, align 4
  invoke void @_ZN7hipsycl2rt5hints14bind_to_deviceC2ENS0_9device_idE(ptr noundef nonnull align 4 dereferenceable(20) %10, i64 %43, i64 %45)
          to label %46 unwind label %126

46:                                               ; preds = %41
  invoke void @_ZN7hipsycl2rt15execution_hints8set_hintINS0_5hints14bind_to_deviceETnNSt9enable_ifIXsr3stdE12is_base_of_vINS3_14execution_hintET_EEiE4typeELi0EEEvOS7_(ptr noundef nonnull align 8 dereferenceable(132) %4, ptr noundef nonnull align 4 dereferenceable(20) %10)
          to label %47 unwind label %126

47:                                               ; preds = %46
  call void @llvm.lifetime.end.p0(i64 20, ptr %10) #8
  call void @llvm.lifetime.start.p0(i64 368, ptr %12) #8
  %48 = invoke noundef ptr @_ZNK7hipsycl2rt24runtime_keep_alive_token3getEv(ptr noundef nonnull align 8 dereferenceable(16) %3)
          to label %49 unwind label %130

49:                                               ; preds = %47
  invoke void @_ZN7hipsycl2rt11dag_builderC1EPNS0_7runtimeE(ptr noundef nonnull align 8 dereferenceable(368) %12, ptr noundef %48)
          to label %50 unwind label %130

50:                                               ; preds = %49
  call void @llvm.lifetime.start.p0(i64 168, ptr %13) #8
  %51 = invoke noundef ptr @_ZNK7hipsycl2rt24runtime_keep_alive_token3getEv(ptr noundef nonnull align 8 dereferenceable(16) %3)
          to label %52 unwind label %134

52:                                               ; preds = %50
  invoke void @_ZN7hipsycl2rt17requirements_listC2EPNS0_7runtimeE(ptr noundef nonnull align 8 dereferenceable(168) %13, ptr noundef %51)
          to label %53 unwind label %134

53:                                               ; preds = %52
  call void @llvm.lifetime.start.p0(i64 8, ptr %14) #8
  call void @llvm.lifetime.start.p0(i64 160, ptr %16) #8
  %54 = getelementptr inbounds %"struct.hipsycl::glue::kernel_launcher_data", ptr %16, i32 0, i32 0
  store ptr null, ptr %54, align 8, !tbaa !18
  %55 = getelementptr inbounds %"struct.hipsycl::glue::kernel_launcher_data", ptr %16, i32 0, i32 1
  store i32 0, ptr %55, align 8, !tbaa !30
  %56 = getelementptr inbounds %"struct.hipsycl::glue::kernel_launcher_data", ptr %16, i32 0, i32 2
  call void @llvm.memset.p0.i64(ptr align 8 %56, i8 0, i64 24, i1 false)
  call void @_ZNSt6vectorIhSaIhEEC2Ev(ptr noundef nonnull align 8 dereferenceable(24) %56) #8
  %57 = getelementptr inbounds %"struct.hipsycl::glue::kernel_launcher_data", ptr %16, i32 0, i32 3
  call void @llvm.memset.p0.i64(ptr align 8 %57, i8 0, i64 24, i1 false)
  %58 = getelementptr inbounds %"struct.hipsycl::glue::kernel_launcher_data", ptr %16, i32 0, i32 4
  call void @llvm.memset.p0.i64(ptr align 8 %58, i8 0, i64 24, i1 false)
  %59 = getelementptr inbounds %"struct.hipsycl::glue::kernel_launcher_data", ptr %16, i32 0, i32 5
  store i32 0, ptr %59, align 8, !tbaa !31
  %60 = getelementptr inbounds %"struct.hipsycl::glue::kernel_launcher_data", ptr %16, i32 0, i32 6
  call void @_ZNSt8functionIFvPN7hipsycl2rt16kernel_operationERNS0_4sycl14interop_handleEEEC2Ev(ptr noundef nonnull align 8 dereferenceable(32) %60) #8
  %61 = getelementptr inbounds %"struct.hipsycl::glue::kernel_launcher_data", ptr %16, i32 0, i32 7
  store i64 0, ptr %61, align 8, !tbaa !32
  %62 = getelementptr inbounds %"struct.hipsycl::glue::kernel_launcher_data", ptr %16, i32 0, i32 8
  store ptr null, ptr %62, align 8, !tbaa !33
  %63 = getelementptr inbounds %"struct.hipsycl::glue::kernel_launcher_data", ptr %16, i32 0, i32 9
  store ptr null, ptr %63, align 8, !tbaa !34
  %64 = getelementptr inbounds %"struct.hipsycl::glue::kernel_launcher_data", ptr %16, i32 0, i32 10
  store ptr null, ptr %64, align 8, !tbaa !35
  call void @_ZN3sbo12small_vectorISt10unique_ptrIN7hipsycl2rt23backend_kernel_launcherESt14default_deleteIS4_EELm8EEC2Ev(ptr noundef nonnull align 8 dereferenceable(96) %17) #8
  invoke void @_ZN7hipsycl2rt15kernel_launcherC2ERKNS_4glue20kernel_launcher_dataEN3sbo12small_vectorISt10unique_ptrINS0_23backend_kernel_launcherESt14default_deleteIS9_EELm8EEE(ptr noundef nonnull align 8 dereferenceable(416) %15, ptr noundef nonnull align 8 dereferenceable(160) %16, ptr noundef %17)
          to label %65 unwind label %138

65:                                               ; preds = %53
  invoke void @_ZN7hipsycl2rt17requirements_listC2ERKS1_(ptr noundef nonnull align 8 dereferenceable(168) %18, ptr noundef nonnull align 8 dereferenceable(168) %13)
          to label %66 unwind label %142

66:                                               ; preds = %65
  invoke void @_ZN7hipsycl2rt14make_operationINS0_16kernel_operationEJPKcNS0_15kernel_launcherENS0_17requirements_listEEEESt10unique_ptrINS0_9operationESt14default_deleteIS8_EEDpT0_(ptr dead_on_unwind writable sret(%"class.std::unique_ptr") align 8 %14, ptr noundef @.str.6, ptr noundef %15, ptr noundef %18)
          to label %67 unwind label %146

67:                                               ; preds = %66
  call void @_ZN7hipsycl2rt17requirements_listD2Ev(ptr noundef nonnull align 8 dereferenceable(168) %18) #8
  call void @_ZN7hipsycl2rt15kernel_launcherD2Ev(ptr noundef nonnull align 8 dereferenceable(416) %15) #8
  call void @_ZNSt6vectorISt10unique_ptrIN7hipsycl2rt23backend_kernel_launcherESt14default_deleteIS3_EEN3sbo29small_buffer_vector_allocatorIS6_Lm8ES6_EEED2Ev(ptr noundef nonnull align 8 dereferenceable(96) %17) #8
  call void @_ZN7hipsycl4glue20kernel_launcher_dataD2Ev(ptr noundef nonnull align 8 dereferenceable(160) %16) #8
  call void @llvm.lifetime.end.p0(i64 160, ptr %16) #8
  call void @llvm.lifetime.start.p0(i64 16, ptr %19) #8
  call void @_ZNSt10unique_ptrIN7hipsycl2rt9operationESt14default_deleteIS2_EEC2EOS5_(ptr noundef nonnull align 8 dereferenceable(8) %20, ptr noundef nonnull align 8 dereferenceable(8) %14) #8
  invoke void @_ZN7hipsycl2rt11dag_builder17add_command_groupESt10unique_ptrINS0_9operationESt14default_deleteIS3_EERKNS0_17requirements_listERKNS0_15execution_hintsE(ptr dead_on_unwind writable sret(%"class.std::shared_ptr.53") align 8 %19, ptr noundef nonnull align 8 dereferenceable(368) %12, ptr noundef %20, ptr noundef nonnull align 8 dereferenceable(168) %13, ptr noundef nonnull align 8 dereferenceable(132) %4)
          to label %68 unwind label %152

68:                                               ; preds = %67
  call void @_ZNSt10unique_ptrIN7hipsycl2rt9operationESt14default_deleteIS2_EED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %20) #8
  call void @llvm.lifetime.start.p0(i64 8, ptr %21) #8
  %69 = call noundef ptr @_ZNKSt19__shared_ptr_accessIN7hipsycl2rt8dag_nodeELN9__gnu_cxx12_Lock_policyE2ELb0ELb0EEptEv(ptr noundef nonnull align 1 dereferenceable(1) %19) #8
  %70 = invoke noundef nonnull align 8 dereferenceable(132) ptr @_ZN7hipsycl2rt8dag_node19get_execution_hintsEv(ptr noundef nonnull align 8 dereferenceable(384) %69)
          to label %71 unwind label %156

71:                                               ; preds = %68
  store ptr %70, ptr %21, align 8, !tbaa !5
  br label %72

72:                                               ; preds = %90, %71
  %73 = load ptr, ptr @_ZN5boost9unit_test12_GLOBAL__N_113unit_test_logE, align 8, !tbaa !5
  invoke void @_ZN5boost9unit_test13basic_cstringIKcEC2ImEEPS2_T_(ptr noundef nonnull align 8 dereferenceable(16) %22, ptr noundef @.str.3, i64 noundef 86)
          to label %74 unwind label %156

74:                                               ; preds = %72
  invoke void @_ZN5boost9unit_test13basic_cstringIKcEC2Ev(ptr noundef nonnull align 8 dereferenceable(16) %23)
          to label %75 unwind label %156

75:                                               ; preds = %74
  invoke void @_ZN5boost9unit_test15unit_test_log_t14set_checkpointENS0_13basic_cstringIKcEEmS4_(ptr noundef nonnull align 8 dereferenceable(8) %73, ptr noundef %22, i64 noundef 48, ptr noundef %23)
          to label %76 unwind label %156

76:                                               ; preds = %75
  call void @llvm.lifetime.start.p0(i64 24, ptr %24) #8
  %77 = load ptr, ptr %21, align 8, !tbaa !5
  %78 = invoke noundef zeroext i1 @_ZNK7hipsycl2rt15execution_hints8has_hintINS0_5hints14bind_to_deviceEEEbv(ptr noundef nonnull align 8 dereferenceable(132) %77)
          to label %79 unwind label %160

79:                                               ; preds = %76
  invoke void @_ZN5boost10test_tools16assertion_resultC2Eb(ptr noundef nonnull align 8 dereferenceable(24) %24, i1 noundef zeroext %78)
          to label %80 unwind label %160

80:                                               ; preds = %79
  call void @llvm.lifetime.start.p0(i64 32, ptr %25) #8
  %81 = invoke noundef nonnull align 8 dereferenceable(9) ptr @_ZN5boost9unit_test12lazy_ostream8instanceEv()
          to label %82 unwind label %164

82:                                               ; preds = %80
  call void @llvm.lifetime.start.p0(i64 16, ptr %26) #8
  invoke void @_ZN5boost9unit_test13basic_cstringIKcEC2ImEEPS2_T_(ptr noundef nonnull align 8 dereferenceable(16) %26, ptr noundef @.str.7, i64 noundef 48)
          to label %83 unwind label %168

83:                                               ; preds = %82
  invoke void @_ZN5boost9unit_testlsINS0_13basic_cstringIKcEEEENS0_17lazy_ostream_implINS0_12lazy_ostreamET_RKS7_EERKS6_S9_(ptr dead_on_unwind writable sret(%"class.boost::unit_test::lazy_ostream_impl") align 8 %25, ptr noundef nonnull align 8 dereferenceable(9) %81, ptr noundef nonnull align 8 dereferenceable(16) %26)
          to label %84 unwind label %168

84:                                               ; preds = %83
  invoke void @_ZN5boost9unit_test13basic_cstringIKcEC2ImEEPS2_T_(ptr noundef nonnull align 8 dereferenceable(16) %27, ptr noundef @.str.3, i64 noundef 86)
          to label %85 unwind label %172

85:                                               ; preds = %84
  %86 = invoke noundef zeroext i1 (ptr, ptr, ptr, i64, i32, i32, i64, ...) @_ZN5boost10test_tools9tt_detail16report_assertionERKNS0_16assertion_resultERKNS_9unit_test12lazy_ostreamENS5_13basic_cstringIKcEEmNS1_10tool_levelENS1_10check_typeEmz(ptr noundef nonnull align 8 dereferenceable(24) %24, ptr noundef nonnull align 8 dereferenceable(9) %25, ptr noundef %27, i64 noundef 48, i32 noundef 1, i32 noundef 0, i64 noundef 0)
          to label %87 unwind label %172

87:                                               ; preds = %85
  call void @_ZN5boost9unit_test12lazy_ostreamD2Ev(ptr noundef nonnull align 8 dereferenceable(32) %25) #8
  call void @llvm.lifetime.end.p0(i64 16, ptr %26) #8
  call void @llvm.lifetime.end.p0(i64 32, ptr %25) #8
  call void @_ZN5boost10test_tools16assertion_resultD2Ev(ptr noundef nonnull align 8 dereferenceable(24) %24) #8
  call void @llvm.lifetime.end.p0(i64 24, ptr %24) #8
  br label %88

88:                                               ; preds = %87
  %89 = invoke noundef zeroext i1 @_ZN5boost10test_tools9tt_detail10dummy_condEv()
          to label %90 unwind label %156

90:                                               ; preds = %88
  br i1 %89, label %72, label %91, !llvm.loop !36

91:                                               ; preds = %90
  br label %92

92:                                               ; preds = %118, %91
  %93 = load ptr, ptr @_ZN5boost9unit_test12_GLOBAL__N_113unit_test_logE, align 8, !tbaa !5
  invoke void @_ZN5boost9unit_test13basic_cstringIKcEC2ImEEPS2_T_(ptr noundef nonnull align 8 dereferenceable(16) %28, ptr noundef @.str.3, i64 noundef 86)
          to label %94 unwind label %156

94:                                               ; preds = %92
  invoke void @_ZN5boost9unit_test13basic_cstringIKcEC2Ev(ptr noundef nonnull align 8 dereferenceable(16) %29)
          to label %95 unwind label %156

95:                                               ; preds = %94
  invoke void @_ZN5boost9unit_test15unit_test_log_t14set_checkpointENS0_13basic_cstringIKcEEmS4_(ptr noundef nonnull align 8 dereferenceable(8) %93, ptr noundef %28, i64 noundef 50, ptr noundef %29)
          to label %96 unwind label %156

96:                                               ; preds = %95
  call void @llvm.lifetime.start.p0(i64 24, ptr %30) #8
  call void @llvm.lifetime.start.p0(i64 16, ptr %31) #8
  %97 = load ptr, ptr %21, align 8, !tbaa !5
  %98 = invoke noundef ptr @_ZNK7hipsycl2rt15execution_hints8get_hintINS0_5hints14bind_to_deviceETnNSt9enable_ifIXsr3stdE12is_base_of_vINS3_14execution_hintET_EEiE4typeELi0EEEPKS7_v(ptr noundef nonnull align 8 dereferenceable(132) %97)
          to label %99 unwind label %179

99:                                               ; preds = %96
  %100 = invoke { i64, i64 } @_ZNK7hipsycl2rt5hints14bind_to_device13get_device_idEv(ptr noundef nonnull align 4 dereferenceable(20) %98)
          to label %101 unwind label %179

101:                                              ; preds = %99
  %102 = getelementptr inbounds { i64, i64 }, ptr %31, i32 0, i32 0
  %103 = extractvalue { i64, i64 } %100, 0
  store i64 %103, ptr %102, align 4
  %104 = getelementptr inbounds { i64, i64 }, ptr %31, i32 0, i32 1
  %105 = extractvalue { i64, i64 } %100, 1
  store i64 %105, ptr %104, align 4
  %106 = invoke noundef zeroext i1 @_ZN7hipsycl2rteqERKNS0_9device_idES3_(ptr noundef nonnull align 4 dereferenceable(16) %31, ptr noundef nonnull align 4 dereferenceable(16) %5)
          to label %107 unwind label %179

107:                                              ; preds = %101
  invoke void @_ZN5boost10test_tools16assertion_resultC2Eb(ptr noundef nonnull align 8 dereferenceable(24) %30, i1 noundef zeroext %106)
          to label %108 unwind label %179

108:                                              ; preds = %107
  call void @llvm.lifetime.start.p0(i64 32, ptr %32) #8
  %109 = invoke noundef nonnull align 8 dereferenceable(9) ptr @_ZN5boost9unit_test12lazy_ostream8instanceEv()
          to label %110 unwind label %183

110:                                              ; preds = %108
  call void @llvm.lifetime.start.p0(i64 16, ptr %33) #8
  invoke void @_ZN5boost9unit_test13basic_cstringIKcEC2ImEEPS2_T_(ptr noundef nonnull align 8 dereferenceable(16) %33, ptr noundef @.str.8, i64 noundef 71)
          to label %111 unwind label %187

111:                                              ; preds = %110
  invoke void @_ZN5boost9unit_testlsINS0_13basic_cstringIKcEEEENS0_17lazy_ostream_implINS0_12lazy_ostreamET_RKS7_EERKS6_S9_(ptr dead_on_unwind writable sret(%"class.boost::unit_test::lazy_ostream_impl") align 8 %32, ptr noundef nonnull align 8 dereferenceable(9) %109, ptr noundef nonnull align 8 dereferenceable(16) %33)
          to label %112 unwind label %187

112:                                              ; preds = %111
  invoke void @_ZN5boost9unit_test13basic_cstringIKcEC2ImEEPS2_T_(ptr noundef nonnull align 8 dereferenceable(16) %34, ptr noundef @.str.3, i64 noundef 86)
          to label %113 unwind label %191

113:                                              ; preds = %112
  %114 = invoke noundef zeroext i1 (ptr, ptr, ptr, i64, i32, i32, i64, ...) @_ZN5boost10test_tools9tt_detail16report_assertionERKNS0_16assertion_resultERKNS_9unit_test12lazy_ostreamENS5_13basic_cstringIKcEEmNS1_10tool_levelENS1_10check_typeEmz(ptr noundef nonnull align 8 dereferenceable(24) %30, ptr noundef nonnull align 8 dereferenceable(9) %32, ptr noundef %34, i64 noundef 50, i32 noundef 1, i32 noundef 0, i64 noundef 0)
          to label %115 unwind label %191

115:                                              ; preds = %113
  call void @_ZN5boost9unit_test12lazy_ostreamD2Ev(ptr noundef nonnull align 8 dereferenceable(32) %32) #8
  call void @llvm.lifetime.end.p0(i64 16, ptr %33) #8
  call void @llvm.lifetime.end.p0(i64 32, ptr %32) #8
  call void @_ZN5boost10test_tools16assertion_resultD2Ev(ptr noundef nonnull align 8 dereferenceable(24) %30) #8
  call void @llvm.lifetime.end.p0(i64 16, ptr %31) #8
  call void @llvm.lifetime.end.p0(i64 24, ptr %30) #8
  br label %116

116:                                              ; preds = %115
  %117 = invoke noundef zeroext i1 @_ZN5boost10test_tools9tt_detail10dummy_condEv()
          to label %118 unwind label %156

118:                                              ; preds = %116
  br i1 %117, label %92, label %119, !llvm.loop !38

119:                                              ; preds = %118
  %120 = call noundef ptr @_ZNKSt19__shared_ptr_accessIN7hipsycl2rt8dag_nodeELN9__gnu_cxx12_Lock_policyE2ELb0ELb0EEptEv(ptr noundef nonnull align 1 dereferenceable(1) %19) #8
  invoke void @_ZN7hipsycl2rt8dag_node6cancelEv(ptr noundef nonnull align 8 dereferenceable(384) %120)
          to label %121 unwind label %156

121:                                              ; preds = %119
  call void @llvm.lifetime.end.p0(i64 8, ptr %21) #8
  call void @_ZNSt12__shared_ptrIN7hipsycl2rt8dag_nodeELN9__gnu_cxx12_Lock_policyE2EED2Ev(ptr noundef nonnull align 8 dereferenceable(16) %19) #8
  call void @llvm.lifetime.end.p0(i64 16, ptr %19) #8
  call void @_ZNSt10unique_ptrIN7hipsycl2rt9operationESt14default_deleteIS2_EED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %14) #8
  call void @llvm.lifetime.end.p0(i64 8, ptr %14) #8
  call void @_ZN7hipsycl2rt17requirements_listD2Ev(ptr noundef nonnull align 8 dereferenceable(168) %13) #8
  call void @llvm.lifetime.end.p0(i64 168, ptr %13) #8
  call void @_ZN7hipsycl2rt11dag_builderD2Ev(ptr noundef nonnull align 8 dereferenceable(368) %12) #8
  call void @llvm.lifetime.end.p0(i64 368, ptr %12) #8
  call void @llvm.lifetime.end.p0(i64 16, ptr %5) #8
  call void @_ZN7hipsycl2rt15execution_hintsD2Ev(ptr noundef nonnull align 8 dereferenceable(132) %4) #8
  call void @llvm.lifetime.end.p0(i64 136, ptr %4) #8
  call void @_ZN7hipsycl2rt24runtime_keep_alive_tokenD2Ev(ptr noundef nonnull align 8 dereferenceable(16) %3) #8
  call void @llvm.lifetime.end.p0(i64 16, ptr %3) #8
  ret void

122:                                              ; preds = %36, %1
  %123 = landingpad { ptr, i32 }
          cleanup
  %124 = extractvalue { ptr, i32 } %123, 0
  store ptr %124, ptr %7, align 8
  %125 = extractvalue { ptr, i32 } %123, 1
  store i32 %125, ptr %8, align 4
  br label %203

126:                                              ; preds = %46, %41
  %127 = landingpad { ptr, i32 }
          cleanup
  %128 = extractvalue { ptr, i32 } %127, 0
  store ptr %128, ptr %7, align 8
  %129 = extractvalue { ptr, i32 } %127, 1
  store i32 %129, ptr %8, align 4
  call void @llvm.lifetime.end.p0(i64 20, ptr %10) #8
  br label %203

130:                                              ; preds = %49, %47
  %131 = landingpad { ptr, i32 }
          cleanup
  %132 = extractvalue { ptr, i32 } %131, 0
  store ptr %132, ptr %7, align 8
  %133 = extractvalue { ptr, i32 } %131, 1
  store i32 %133, ptr %8, align 4
  br label %202

134:                                              ; preds = %52, %50
  %135 = landingpad { ptr, i32 }
          cleanup
  %136 = extractvalue { ptr, i32 } %135, 0
  store ptr %136, ptr %7, align 8
  %137 = extractvalue { ptr, i32 } %135, 1
  store i32 %137, ptr %8, align 4
  br label %201

138:                                              ; preds = %53
  %139 = landingpad { ptr, i32 }
          cleanup
  %140 = extractvalue { ptr, i32 } %139, 0
  store ptr %140, ptr %7, align 8
  %141 = extractvalue { ptr, i32 } %139, 1
  store i32 %141, ptr %8, align 4
  br label %151

142:                                              ; preds = %65
  %143 = landingpad { ptr, i32 }
          cleanup
  %144 = extractvalue { ptr, i32 } %143, 0
  store ptr %144, ptr %7, align 8
  %145 = extractvalue { ptr, i32 } %143, 1
  store i32 %145, ptr %8, align 4
  br label %150

146:                                              ; preds = %66
  %147 = landingpad { ptr, i32 }
          cleanup
  %148 = extractvalue { ptr, i32 } %147, 0
  store ptr %148, ptr %7, align 8
  %149 = extractvalue { ptr, i32 } %147, 1
  store i32 %149, ptr %8, align 4
  call void @_ZN7hipsycl2rt17requirements_listD2Ev(ptr noundef nonnull align 8 dereferenceable(168) %18) #8
  br label %150

150:                                              ; preds = %146, %142
  call void @_ZN7hipsycl2rt15kernel_launcherD2Ev(ptr noundef nonnull align 8 dereferenceable(416) %15) #8
  br label %151

151:                                              ; preds = %150, %138
  call void @_ZNSt6vectorISt10unique_ptrIN7hipsycl2rt23backend_kernel_launcherESt14default_deleteIS3_EEN3sbo29small_buffer_vector_allocatorIS6_Lm8ES6_EEED2Ev(ptr noundef nonnull align 8 dereferenceable(96) %17) #8
  call void @_ZN7hipsycl4glue20kernel_launcher_dataD2Ev(ptr noundef nonnull align 8 dereferenceable(160) %16) #8
  call void @llvm.lifetime.end.p0(i64 160, ptr %16) #8
  br label %200

152:                                              ; preds = %67
  %153 = landingpad { ptr, i32 }
          cleanup
  %154 = extractvalue { ptr, i32 } %153, 0
  store ptr %154, ptr %7, align 8
  %155 = extractvalue { ptr, i32 } %153, 1
  store i32 %155, ptr %8, align 4
  call void @_ZNSt10unique_ptrIN7hipsycl2rt9operationESt14default_deleteIS2_EED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %20) #8
  br label %199

156:                                              ; preds = %119, %116, %95, %94, %92, %88, %75, %74, %72, %68
  %157 = landingpad { ptr, i32 }
          cleanup
  %158 = extractvalue { ptr, i32 } %157, 0
  store ptr %158, ptr %7, align 8
  %159 = extractvalue { ptr, i32 } %157, 1
  store i32 %159, ptr %8, align 4
  br label %198

160:                                              ; preds = %79, %76
  %161 = landingpad { ptr, i32 }
          cleanup
  %162 = extractvalue { ptr, i32 } %161, 0
  store ptr %162, ptr %7, align 8
  %163 = extractvalue { ptr, i32 } %161, 1
  store i32 %163, ptr %8, align 4
  br label %178

164:                                              ; preds = %80
  %165 = landingpad { ptr, i32 }
          cleanup
  %166 = extractvalue { ptr, i32 } %165, 0
  store ptr %166, ptr %7, align 8
  %167 = extractvalue { ptr, i32 } %165, 1
  store i32 %167, ptr %8, align 4
  br label %177

168:                                              ; preds = %83, %82
  %169 = landingpad { ptr, i32 }
          cleanup
  %170 = extractvalue { ptr, i32 } %169, 0
  store ptr %170, ptr %7, align 8
  %171 = extractvalue { ptr, i32 } %169, 1
  store i32 %171, ptr %8, align 4
  br label %176

172:                                              ; preds = %85, %84
  %173 = landingpad { ptr, i32 }
          cleanup
  %174 = extractvalue { ptr, i32 } %173, 0
  store ptr %174, ptr %7, align 8
  %175 = extractvalue { ptr, i32 } %173, 1
  store i32 %175, ptr %8, align 4
  call void @_ZN5boost9unit_test12lazy_ostreamD2Ev(ptr noundef nonnull align 8 dereferenceable(32) %25) #8
  br label %176

176:                                              ; preds = %172, %168
  call void @llvm.lifetime.end.p0(i64 16, ptr %26) #8
  br label %177

177:                                              ; preds = %176, %164
  call void @llvm.lifetime.end.p0(i64 32, ptr %25) #8
  call void @_ZN5boost10test_tools16assertion_resultD2Ev(ptr noundef nonnull align 8 dereferenceable(24) %24) #8
  br label %178

178:                                              ; preds = %177, %160
  call void @llvm.lifetime.end.p0(i64 24, ptr %24) #8
  br label %198

179:                                              ; preds = %107, %101, %99, %96
  %180 = landingpad { ptr, i32 }
          cleanup
  %181 = extractvalue { ptr, i32 } %180, 0
  store ptr %181, ptr %7, align 8
  %182 = extractvalue { ptr, i32 } %180, 1
  store i32 %182, ptr %8, align 4
  br label %197

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

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

191:                                              ; preds = %113, %112
  %192 = landingpad { ptr, i32 }
          cleanup
  %193 = extractvalue { ptr, i32 } %192, 0
  store ptr %193, ptr %7, align 8
  %194 = extractvalue { ptr, i32 } %192, 1
  store i32 %194, ptr %8, align 4
  call void @_ZN5boost9unit_test12lazy_ostreamD2Ev(ptr noundef nonnull align 8 dereferenceable(32) %32) #8
  br label %195

195:                                              ; preds = %191, %187
  call void @llvm.lifetime.end.p0(i64 16, ptr %33) #8
  br label %196

196:                                              ; preds = %195, %183
  call void @llvm.lifetime.end.p0(i64 32, ptr %32) #8
  call void @_ZN5boost10test_tools16assertion_resultD2Ev(ptr noundef nonnull align 8 dereferenceable(24) %30) #8
  br label %197

197:                                              ; preds = %196, %179
  call void @llvm.lifetime.end.p0(i64 16, ptr %31) #8
  call void @llvm.lifetime.end.p0(i64 24, ptr %30) #8
  br label %198

198:                                              ; preds = %197, %178, %156
  call void @llvm.lifetime.end.p0(i64 8, ptr %21) #8
  call void @_ZNSt12__shared_ptrIN7hipsycl2rt8dag_nodeELN9__gnu_cxx12_Lock_policyE2EED2Ev(ptr noundef nonnull align 8 dereferenceable(16) %19) #8
  br label %199

199:                                              ; preds = %198, %152
  call void @llvm.lifetime.end.p0(i64 16, ptr %19) #8
  call void @_ZNSt10unique_ptrIN7hipsycl2rt9operationESt14default_deleteIS2_EED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %14) #8
  br label %200

200:                                              ; preds = %199, %151
  call void @llvm.lifetime.end.p0(i64 8, ptr %14) #8
  call void @_ZN7hipsycl2rt17requirements_listD2Ev(ptr noundef nonnull align 8 dereferenceable(168) %13) #8
  br label %201

201:                                              ; preds = %200, %134
  call void @llvm.lifetime.end.p0(i64 168, ptr %13) #8
  call void @_ZN7hipsycl2rt11dag_builderD2Ev(ptr noundef nonnull align 8 dereferenceable(368) %12) #8
  br label %202

202:                                              ; preds = %201, %130
  call void @llvm.lifetime.end.p0(i64 368, ptr %12) #8
  br label %203

203:                                              ; preds = %202, %126, %122
  call void @llvm.lifetime.end.p0(i64 16, ptr %5) #8
  call void @_ZN7hipsycl2rt15execution_hintsD2Ev(ptr noundef nonnull align 8 dereferenceable(132) %4) #8
  call void @llvm.lifetime.end.p0(i64 136, ptr %4) #8
  call void @_ZN7hipsycl2rt24runtime_keep_alive_tokenD2Ev(ptr noundef nonnull align 8 dereferenceable(16) %3) #8
  call void @llvm.lifetime.end.p0(i64 16, ptr %3) #8
  br label %204

204:                                              ; preds = %203
  %205 = load ptr, ptr %7, align 8
  %206 = load i32, ptr %8, align 4
  %207 = insertvalue { ptr, i32 } poison, ptr %205, 0
  %208 = insertvalue { ptr, i32 } %207, i32 %206, 1
  resume { ptr, i32 } %208
}

declare void @_ZN7hipsycl2rt24runtime_keep_alive_tokenC1Ev(ptr noundef nonnull align 8 dereferenceable(16)) unnamed_addr #2

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

; Function Attrs: mustprogress nounwind uwtable
declare dso_local void @_ZN7hipsycl2rt18backend_descriptorC2ENS0_17hardware_platformENS0_12api_platformE(ptr noundef nonnull align 4 dereferenceable(12), i32 noundef, i32 noundef) unnamed_addr #3 align 2

declare void @_ZN7hipsycl2rt9device_idC1ENS0_18backend_descriptorEi(ptr noundef nonnull align 4 dereferenceable(16), i64, i32, i32 noundef) unnamed_addr #2

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

; Function Attrs: mustprogress uwtable
declare dso_local void @_ZN7hipsycl2rt15execution_hints8set_hintINS0_5hints14bind_to_deviceETnNSt9enable_ifIXsr3stdE12is_base_of_vINS3_14execution_hintET_EEiE4typeELi0EEEvOS7_(ptr noundef nonnull align 8 dereferenceable(132), ptr noundef nonnull align 4 dereferenceable(20)) #1 align 2

; Function Attrs: mustprogress nounwind uwtable
declare dso_local void @_ZN7hipsycl2rt5hints14bind_to_deviceC2ENS0_9device_idE(ptr noundef nonnull align 4 dereferenceable(20), i64, i64) unnamed_addr #3 align 2

declare noundef ptr @_ZNK7hipsycl2rt24runtime_keep_alive_token3getEv(ptr noundef nonnull align 8 dereferenceable(16)) #2

declare void @_ZN7hipsycl2rt11dag_builderC1EPNS0_7runtimeE(ptr noundef nonnull align 8 dereferenceable(368), ptr noundef) unnamed_addr #2

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

; Function Attrs: mustprogress uwtable
declare dso_local void @_ZN7hipsycl2rt14make_operationINS0_16kernel_operationEJPKcNS0_15kernel_launcherENS0_17requirements_listEEEESt10unique_ptrINS0_9operationESt14default_deleteIS8_EEDpT0_(ptr dead_on_unwind noalias writable sret(%"class.std::unique_ptr") align 8, ptr noundef, ptr noundef, ptr noundef) #1

; Function Attrs: nocallback nofree nounwind willreturn memory(argmem: write)
declare void @llvm.memset.p0.i64(ptr nocapture writeonly, i8, i64, i1 immarg) #5

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

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

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

; Function Attrs: mustprogress uwtable
declare dso_local void @_ZN7hipsycl2rt15kernel_launcherC2ERKNS_4glue20kernel_launcher_dataEN3sbo12small_vectorISt10unique_ptrINS0_23backend_kernel_launcherESt14default_deleteIS9_EELm8EEE(ptr noundef nonnull align 8 dereferenceable(416), ptr noundef nonnull align 8 dereferenceable(160), ptr noundef) unnamed_addr #1 align 2

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

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

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

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

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

declare void @_ZN7hipsycl2rt11dag_builder17add_command_groupESt10unique_ptrINS0_9operationESt14default_deleteIS3_EERKNS0_17requirements_listERKNS0_15execution_hintsE(ptr dead_on_unwind writable sret(%"class.std::shared_ptr.53") align 8, ptr noundef nonnull align 8 dereferenceable(368), ptr noundef, ptr noundef nonnull align 8 dereferenceable(168), ptr noundef nonnull align 8 dereferenceable(132)) #2

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

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

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

declare noundef nonnull align 8 dereferenceable(132) ptr @_ZN7hipsycl2rt8dag_node19get_execution_hintsEv(ptr noundef nonnull align 8 dereferenceable(384)) #2

declare void @_ZN5boost9unit_test15unit_test_log_t14set_checkpointENS0_13basic_cstringIKcEEmS4_(ptr noundef nonnull align 8 dereferenceable(8), ptr noundef, i64 noundef, ptr noundef) #2

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

; Function Attrs: inlinehint mustprogress uwtable
declare dso_local void @_ZN5boost9unit_test13basic_cstringIKcEC2Ev(ptr noundef nonnull align 8 dereferenceable(16)) unnamed_addr #6 align 2

declare noundef zeroext i1 @_ZN5boost10test_tools9tt_detail16report_assertionERKNS0_16assertion_resultERKNS_9unit_test12lazy_ostreamENS5_13basic_cstringIKcEEmNS1_10tool_levelENS1_10check_typeEmz(ptr noundef nonnull align 8 dereferenceable(24), ptr noundef nonnull align 8 dereferenceable(9), ptr noundef, i64 noundef, i32 noundef, i32 noundef, i64 noundef, ...) #2

; Function Attrs: mustprogress uwtable
declare dso_local noundef zeroext i1 @_ZNK7hipsycl2rt15execution_hints8has_hintINS0_5hints14bind_to_deviceEEEbv(ptr noundef nonnull align 8 dereferenceable(132)) #1 align 2

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

; Function Attrs: inlinehint mustprogress uwtable
declare dso_local void @_ZN5boost9unit_testlsINS0_13basic_cstringIKcEEEENS0_17lazy_ostream_implINS0_12lazy_ostreamET_RKS7_EERKS6_S9_(ptr dead_on_unwind noalias writable sret(%"class.boost::unit_test::lazy_ostream_impl") align 8, ptr noundef nonnull align 8 dereferenceable(9), ptr noundef nonnull align 8 dereferenceable(16)) #6

; Function Attrs: mustprogress nounwind uwtable
declare dso_local noundef nonnull align 8 dereferenceable(9) ptr @_ZN5boost9unit_test12lazy_ostream8instanceEv() #3 align 2

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

; Function Attrs: inlinehint mustprogress nounwind uwtable
declare dso_local noundef zeroext i1 @_ZN5boost10test_tools9tt_detail10dummy_condEv() #7

; Function Attrs: mustprogress uwtable
declare dso_local noundef zeroext i1 @_ZN7hipsycl2rteqERKNS0_9device_idES3_(ptr noundef nonnull align 4 dereferenceable(16), ptr noundef nonnull align 4 dereferenceable(16)) #1

; Function Attrs: mustprogress nounwind uwtable
declare dso_local noundef ptr @_ZNK7hipsycl2rt15execution_hints8get_hintINS0_5hints14bind_to_deviceETnNSt9enable_ifIXsr3stdE12is_base_of_vINS3_14execution_hintET_EEiE4typeELi0EEEPKS7_v(ptr noundef nonnull align 8 dereferenceable(132)) #3 align 2

; Function Attrs: mustprogress nounwind uwtable
declare dso_local { i64, i64 } @_ZNK7hipsycl2rt5hints14bind_to_device13get_device_idEv(ptr noundef nonnull align 4 dereferenceable(20)) #3 align 2

declare void @_ZN7hipsycl2rt8dag_node6cancelEv(ptr noundef nonnull align 8 dereferenceable(384)) #2

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

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

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

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

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

attributes #0 = { nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) }
attributes #1 = { 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 #2 = { "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cmov,+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #3 = { mustprogress 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 = { nocallback nofree nounwind willreturn memory(argmem: readwrite) }
attributes #5 = { nocallback nofree nounwind willreturn memory(argmem: write) }
attributes #6 = { 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 #7 = { 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 #8 = { 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 = !{i64 0, i64 4, !10, i64 4, i64 4, !12, i64 8, i64 4, !14, i64 12, i64 4, !16}
!10 = !{!11, !11, i64 0}
!11 = !{!"_ZTSN7hipsycl2rt10backend_idE", !7, i64 0}
!12 = !{!13, !13, i64 0}
!13 = !{!"_ZTSN7hipsycl2rt17hardware_platformE", !7, i64 0}
!14 = !{!15, !15, i64 0}
!15 = !{!"_ZTSN7hipsycl2rt12api_platformE", !7, i64 0}
!16 = !{!17, !17, i64 0}
!17 = !{!"int", !7, i64 0}
!18 = !{!19, !6, i64 0}
!19 = !{!"_ZTSN7hipsycl4glue20kernel_launcher_dataE", !6, i64 0, !20, i64 8, !21, i64 16, !25, i64 40, !25, i64 64, !17, i64 88, !27, i64 96, !29, i64 128, !6, i64 136, !6, i64 144, !6, i64 152}
!20 = !{!"_ZTSN7hipsycl2rt11kernel_typeE", !7, i64 0}
!21 = !{!"_ZTSSt6vectorIhSaIhEE", !22, i64 0}
!22 = !{!"_ZTSSt12_Vector_baseIhSaIhEE", !23, i64 0}
!23 = !{!"_ZTSNSt12_Vector_baseIhSaIhEE12_Vector_implE", !24, i64 0}
!24 = !{!"_ZTSNSt12_Vector_baseIhSaIhEE17_Vector_impl_dataE", !6, i64 0, !6, i64 8, !6, i64 16}
!25 = !{!"_ZTSN7hipsycl2rt12static_arrayILi3EEE", !26, i64 0}
!26 = !{!"_ZTSSt5arrayImLm3EE", !7, i64 0}
!27 = !{!"_ZTSSt8functionIFvPN7hipsycl2rt16kernel_operationERNS0_4sycl14interop_handleEEE", !28, i64 0, !6, i64 24}
!28 = !{!"_ZTSSt14_Function_base", !7, i64 0, !6, i64 16}
!29 = !{!"long long", !7, i64 0}
!30 = !{!19, !20, i64 8}
!31 = !{!19, !17, i64 88}
!32 = !{!19, !29, i64 128}
!33 = !{!19, !6, i64 136}
!34 = !{!19, !6, i64 144}
!35 = !{!19, !6, i64 152}
!36 = distinct !{!36, !37}
!37 = !{!"llvm.loop.mustprogress"}
!38 = distinct !{!38, !37}


; 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::runtime_keep_alive_token" = type { %"class.std::shared_ptr" }
%"class.std::shared_ptr" = type { %"class.std::__shared_ptr" }
%"class.std::__shared_ptr" = type { ptr, %"class.std::__shared_count" }
%"class.std::__shared_count" = type { ptr }
%"class.hipsycl::rt::execution_hints" = type <{ %"class.hipsycl::rt::hints::bind_to_device", [4 x i8], %"class.hipsycl::rt::hints::bind_to_device_group", %"class.hipsycl::rt::hints::prefer_execution_lane", %"class.hipsycl::rt::hints::node_group", %"class.hipsycl::rt::hints::coarse_grained_synchronization", [7 x i8], %"class.hipsycl::rt::hints::prefer_executor", %"class.hipsycl::rt::hints::request_instrumentation_submission_timestamp", %"class.hipsycl::rt::hints::request_instrumentation_start_timestamp", %"class.hipsycl::rt::hints::request_instrumentation_finish_timestamp", %"class.hipsycl::rt::hints::instant_execution", [4 x i8] }>
%"class.hipsycl::rt::hints::bind_to_device" = type { %"class.hipsycl::rt::hints::execution_hint", %"class.hipsycl::rt::device_id" }
%"class.hipsycl::rt::hints::execution_hint" = type { i8 }
%"class.hipsycl::rt::device_id" = type { %"struct.hipsycl::rt::backend_descriptor", i32 }
%"struct.hipsycl::rt::backend_descriptor" = type { i32, i32, i32 }
%"class.hipsycl::rt::hints::bind_to_device_group" = type { %"class.hipsycl::rt::hints::execution_hint", %"class.std::vector.0" }
%"class.std::vector.0" = type { %"struct.std::_Vector_base.1" }
%"struct.std::_Vector_base.1" = type { %"struct.std::_Vector_base<hipsycl::rt::device_id, std::allocator<hipsycl::rt::device_id>>::_Vector_impl" }
%"struct.std::_Vector_base<hipsycl::rt::device_id, std::allocator<hipsycl::rt::device_id>>::_Vector_impl" = type { %"struct.std::_Vector_base<hipsycl::rt::device_id, std::allocator<hipsycl::rt::device_id>>::_Vector_impl_data" }
%"struct.std::_Vector_base<hipsycl::rt::device_id, std::allocator<hipsycl::rt::device_id>>::_Vector_impl_data" = type { ptr, ptr, ptr }
%"class.hipsycl::rt::hints::prefer_execution_lane" = type { %"class.hipsycl::rt::hints::execution_hint", i64 }
%"class.hipsycl::rt::hints::node_group" = type { %"class.hipsycl::rt::hints::execution_hint", i64 }
%"class.hipsycl::rt::hints::coarse_grained_synchronization" = type { %"class.hipsycl::rt::hints::execution_hint" }
%"class.hipsycl::rt::hints::prefer_executor" = type { %"class.hipsycl::rt::hints::execution_hint", ptr, %"class.std::shared_ptr.5" }
%"class.std::shared_ptr.5" = type { %"class.std::__shared_ptr.6" }
%"class.std::__shared_ptr.6" = type { ptr, %"class.std::__shared_count" }
%"class.hipsycl::rt::hints::request_instrumentation_submission_timestamp" = type { %"class.hipsycl::rt::hints::execution_hint" }
%"class.hipsycl::rt::hints::request_instrumentation_start_timestamp" = type { %"class.hipsycl::rt::hints::execution_hint" }
%"class.hipsycl::rt::hints::request_instrumentation_finish_timestamp" = type { %"class.hipsycl::rt::hints::execution_hint" }
%"class.hipsycl::rt::hints::instant_execution" = type { %"class.hipsycl::rt::hints::execution_hint" }
%"class.hipsycl::rt::dag_builder" = type { %"class.std::mutex", %"class.hipsycl::rt::dag", ptr }
%"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.hipsycl::rt::dag" = type { %"class.sbo::small_vector", %"class.sbo::small_vector" }
%"class.sbo::small_vector" = type { %"class.std::vector.8" }
%"class.std::vector.8" = type { %"struct.std::_Vector_base.9" }
%"struct.std::_Vector_base.9" = type { %"struct.std::_Vector_base<std::shared_ptr<hipsycl::rt::dag_node>, sbo::small_buffer_vector_allocator<std::shared_ptr<hipsycl::rt::dag_node>>>::_Vector_impl" }
%"struct.std::_Vector_base<std::shared_ptr<hipsycl::rt::dag_node>, sbo::small_buffer_vector_allocator<std::shared_ptr<hipsycl::rt::dag_node>>>::_Vector_impl" = type { %"struct.sbo::small_buffer_vector_allocator.base", %"struct.std::_Vector_base<std::shared_ptr<hipsycl::rt::dag_node>, sbo::small_buffer_vector_allocator<std::shared_ptr<hipsycl::rt::dag_node>>>::_Vector_impl_data" }
%"struct.sbo::small_buffer_vector_allocator.base" = type { [128 x i8], i8, i8 }
%"struct.std::_Vector_base<std::shared_ptr<hipsycl::rt::dag_node>, sbo::small_buffer_vector_allocator<std::shared_ptr<hipsycl::rt::dag_node>>>::_Vector_impl_data" = type { ptr, ptr, ptr }
%"class.hipsycl::rt::requirements_list" = type { %"class.sbo::small_vector", 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.11" }
%"struct.std::_Head_base.11" = type { ptr }
%"class.hipsycl::rt::kernel_launcher" = type { %"class.sbo::small_vector.12", %"struct.hipsycl::glue::kernel_launcher_data", %"class.hipsycl::rt::kernel_configuration" }
%"class.sbo::small_vector.12" = type { %"class.std::vector.13" }
%"class.std::vector.13" = type { %"struct.std::_Vector_base.14" }
%"struct.std::_Vector_base.14" = type { %"struct.std::_Vector_base<std::unique_ptr<hipsycl::rt::backend_kernel_launcher>, sbo::small_buffer_vector_allocator<std::unique_ptr<hipsycl::rt::backend_kernel_launcher>>>::_Vector_impl" }
%"struct.std::_Vector_base<std::unique_ptr<hipsycl::rt::backend_kernel_launcher>, sbo::small_buffer_vector_allocator<std::unique_ptr<hipsycl::rt::backend_kernel_launcher>>>::_Vector_impl" = type { %"struct.sbo::small_buffer_vector_allocator.base.16", %"struct.std::_Vector_base<std::unique_ptr<hipsycl::rt::backend_kernel_launcher>, sbo::small_buffer_vector_allocator<std::unique_ptr<hipsycl::rt::backend_kernel_launcher>>>::_Vector_impl_data" }
%"struct.sbo::small_buffer_vector_allocator.base.16" = type { [64 x i8], i8, i8 }
%"struct.std::_Vector_base<std::unique_ptr<hipsycl::rt::backend_kernel_launcher>, sbo::small_buffer_vector_allocator<std::unique_ptr<hipsycl::rt::backend_kernel_launcher>>>::_Vector_impl_data" = type { ptr, ptr, ptr }
%"struct.hipsycl::glue::kernel_launcher_data" = type { ptr, i32, %"class.std::vector.17", %"class.hipsycl::rt::static_array", %"class.hipsycl::rt::static_array", i32, %"class.std::function", i64, ptr, ptr, ptr }
%"class.std::vector.17" = type { %"struct.std::_Vector_base.18" }
%"struct.std::_Vector_base.18" = type { %"struct.std::_Vector_base<unsigned char, std::allocator<unsigned char>>::_Vector_impl" }
%"struct.std::_Vector_base<unsigned char, std::allocator<unsigned char>>::_Vector_impl" = type { %"struct.std::_Vector_base<unsigned char, std::allocator<unsigned char>>::_Vector_impl_data" }
%"struct.std::_Vector_base<unsigned char, std::allocator<unsigned char>>::_Vector_impl_data" = type { ptr, ptr, ptr }
%"class.hipsycl::rt::static_array" = type { %"struct.std::array" }
%"struct.std::array" = type { [3 x i64] }
%"class.std::function" = type { %"class.std::_Function_base", ptr }
%"class.std::_Function_base" = type { %"union.std::_Any_data", ptr }
%"union.std::_Any_data" = type { %"union.std::_Nocopy_types" }
%"union.std::_Nocopy_types" = type { { i64, i64 } }
%"class.hipsycl::rt::kernel_configuration" = type { %"class.std::vector.22", %"class.std::vector.27", %"class.std::vector.32", %"class.std::vector.37", %"class.std::vector.42", %"class.std::vector.47", %"struct.std::array.52" }
%"class.std::vector.22" = type { %"struct.std::_Vector_base.23" }
%"struct.std::_Vector_base.23" = 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.27" = type { %"struct.std::_Vector_base.28" }
%"struct.std::_Vector_base.28" = 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.32" = type { %"struct.std::_Vector_base.33" }
%"struct.std::_Vector_base.33" = 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.37" = type { %"struct.std::_Vector_base.38" }
%"struct.std::_Vector_base.38" = 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.42" = type { %"struct.std::_Vector_base.43" }
%"struct.std::_Vector_base.43" = 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.std::vector.47" = type { %"struct.std::_Vector_base.48" }
%"struct.std::_Vector_base.48" = 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.52" = type { [2 x i64] }
%"class.std::shared_ptr.53" = type { %"class.std::__shared_ptr.54" }
%"class.std::__shared_ptr.54" = type { ptr, %"class.std::__shared_count" }
%"class.boost::unit_test::basic_cstring" = type { ptr, ptr }
%"class.boost::test_tools::assertion_result" = type { %"class.boost::test_tools::assertion_result::readonly_property65", %"class.boost::shared_ptr" }
%"class.boost::test_tools::assertion_result::readonly_property65" = type { %"class.boost::unit_test::readonly_property" }
%"class.boost::unit_test::readonly_property" = type { %"class.boost::unit_test::class_property" }
%"class.boost::unit_test::class_property" = type { i8 }
%"class.boost::shared_ptr" = type { ptr, %"class.boost::detail::shared_count" }
%"class.boost::detail::shared_count" = type { ptr }
%"class.boost::unit_test::lazy_ostream_impl" = type { %"class.boost::unit_test::lazy_ostream.base", ptr, ptr }
%"class.boost::unit_test::lazy_ostream.base" = type <{ ptr, i8 }>

@_ZN5boost9unit_test12_GLOBAL__N_113unit_test_logE = external hidden local_unnamed_addr global ptr, align 8
@.str.3 = external hidden unnamed_addr constant [87 x i8], align 1
@.str.6 = external hidden unnamed_addr constant [12 x i8], align 1
@.str.7 = external hidden unnamed_addr constant [49 x i8], align 1
@.str.8 = external hidden unnamed_addr constant [72 x i8], align 1

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

declare i32 @__gxx_personality_v0(...)

; 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 uwtable
define dso_local void @_ZN11dag_builder13default_hints11test_methodEv(ptr nocapture noundef nonnull readnone align 1 dereferenceable(1) %0) local_unnamed_addr #1 align 2 personality ptr @__gxx_personality_v0 {
  %2 = alloca %"class.hipsycl::rt::runtime_keep_alive_token", align 8
  %3 = alloca %"class.hipsycl::rt::execution_hints", align 8
  %4 = alloca %"class.hipsycl::rt::device_id", align 8
  %5 = alloca %"struct.hipsycl::rt::backend_descriptor", align 8
  %6 = alloca %"class.hipsycl::rt::hints::bind_to_device", align 4
  %7 = alloca %"class.hipsycl::rt::dag_builder", align 8
  %8 = alloca %"class.hipsycl::rt::requirements_list", align 8
  %9 = alloca %"class.std::unique_ptr", align 8
  %10 = alloca %"class.hipsycl::rt::kernel_launcher", align 8
  %11 = alloca %"struct.hipsycl::glue::kernel_launcher_data", align 8
  %12 = alloca %"class.sbo::small_vector.12", align 8
  %13 = alloca %"class.hipsycl::rt::requirements_list", align 8
  %14 = alloca %"class.std::shared_ptr.53", align 8
  %15 = alloca %"class.std::unique_ptr", align 8
  %16 = alloca %"class.boost::unit_test::basic_cstring", align 8
  %17 = alloca %"class.boost::unit_test::basic_cstring", align 8
  %18 = alloca %"class.boost::test_tools::assertion_result", align 8
  %19 = alloca %"class.boost::unit_test::lazy_ostream_impl", align 8
  %20 = alloca %"class.boost::unit_test::basic_cstring", align 8
  %21 = alloca %"class.boost::unit_test::basic_cstring", align 8
  %22 = alloca %"class.boost::unit_test::basic_cstring", align 8
  %23 = alloca %"class.boost::unit_test::basic_cstring", align 8
  %24 = alloca %"class.boost::test_tools::assertion_result", align 8
  %25 = alloca %"class.hipsycl::rt::device_id", align 8
  %26 = alloca %"class.boost::unit_test::lazy_ostream_impl", align 8
  %27 = alloca %"class.boost::unit_test::basic_cstring", align 8
  %28 = alloca %"class.boost::unit_test::basic_cstring", align 8
  call void @llvm.lifetime.start.p0(i64 16, ptr nonnull %2) #7
  call void @_ZN7hipsycl2rt24runtime_keep_alive_tokenC1Ev(ptr noundef nonnull align 8 dereferenceable(16) %2)
  call void @llvm.lifetime.start.p0(i64 136, ptr nonnull %3) #7
  call void @_ZN7hipsycl2rt15execution_hintsC2Ev(ptr noundef nonnull align 8 dereferenceable(132) %3) #7
  call void @llvm.lifetime.start.p0(i64 16, ptr nonnull %4) #7
  call void @_ZN7hipsycl2rt18backend_descriptorC2ENS0_17hardware_platformENS0_12api_platformE(ptr noundef nonnull align 4 dereferenceable(12) %5, i32 noundef 4, i32 noundef 4)
  %.sroa.02.0.copyload = load i64, ptr %5, align 8
  %.sroa.23.0..sroa_idx = getelementptr inbounds i8, ptr %5, i64 8
  %.sroa.23.0.copyload = load i32, ptr %.sroa.23.0..sroa_idx, align 8
  invoke void @_ZN7hipsycl2rt9device_idC1ENS0_18backend_descriptorEi(ptr noundef nonnull align 4 dereferenceable(16) %4, i64 %.sroa.02.0.copyload, i32 %.sroa.23.0.copyload, i32 noundef 12345)
          to label %29 unwind label %78

29:                                               ; preds = %1
  call void @llvm.lifetime.start.p0(i64 20, ptr nonnull %6) #7
  %.sroa.0.0.copyload = load i64, ptr %4, align 8
  %.sroa.2.0..sroa_idx = getelementptr inbounds i8, ptr %4, i64 8
  %.sroa.2.0.copyload = load i64, ptr %.sroa.2.0..sroa_idx, align 8
  call void @_ZN7hipsycl2rt5hints14bind_to_deviceC2ENS0_9device_idE(ptr noundef nonnull align 4 dereferenceable(20) %6, i64 %.sroa.0.0.copyload, i64 %.sroa.2.0.copyload)
  invoke void @_ZN7hipsycl2rt15execution_hints8set_hintINS0_5hints14bind_to_deviceETnNSt9enable_ifIXsr3stdE12is_base_of_vINS3_14execution_hintET_EEiE4typeELi0EEEvOS7_(ptr noundef nonnull align 8 dereferenceable(132) %3, ptr noundef nonnull align 4 dereferenceable(20) %6)
          to label %30 unwind label %80

30:                                               ; preds = %29
  call void @llvm.lifetime.end.p0(i64 20, ptr nonnull %6) #7
  call void @llvm.lifetime.start.p0(i64 368, ptr nonnull %7) #7
  %31 = invoke noundef ptr @_ZNK7hipsycl2rt24runtime_keep_alive_token3getEv(ptr noundef nonnull align 8 dereferenceable(16) %2)
          to label %32 unwind label %82

32:                                               ; preds = %30
  invoke void @_ZN7hipsycl2rt11dag_builderC1EPNS0_7runtimeE(ptr noundef nonnull align 8 dereferenceable(368) %7, ptr noundef %31)
          to label %33 unwind label %82

33:                                               ; preds = %32
  call void @llvm.lifetime.start.p0(i64 168, ptr nonnull %8) #7
  %34 = invoke noundef ptr @_ZNK7hipsycl2rt24runtime_keep_alive_token3getEv(ptr noundef nonnull align 8 dereferenceable(16) %2)
          to label %35 unwind label %84

35:                                               ; preds = %33
  call void @_ZN7hipsycl2rt17requirements_listC2EPNS0_7runtimeE(ptr noundef nonnull align 8 dereferenceable(168) %8, ptr noundef %34)
  call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %9) #7
  call void @llvm.lifetime.start.p0(i64 160, ptr nonnull %11) #7
  store ptr null, ptr %11, align 8, !tbaa !5
  %36 = getelementptr inbounds i8, ptr %11, i64 8
  store i32 0, ptr %36, align 8, !tbaa !21
  %37 = getelementptr inbounds i8, ptr %11, i64 16
  call void @llvm.memset.p0.i64(ptr noundef nonnull align 8 dereferenceable(24) %37, i8 0, i64 24, i1 false)
  call void @_ZNSt6vectorIhSaIhEEC2Ev(ptr noundef nonnull align 8 dereferenceable(24) %37) #7
  %38 = getelementptr inbounds i8, ptr %11, i64 40
  %39 = getelementptr inbounds i8, ptr %11, i64 96
  call void @llvm.memset.p0.i64(ptr noundef nonnull align 8 dereferenceable(52) %38, i8 0, i64 52, i1 false)
  call void @_ZNSt8functionIFvPN7hipsycl2rt16kernel_operationERNS0_4sycl14interop_handleEEEC2Ev(ptr noundef nonnull align 8 dereferenceable(32) %39) #7
  %40 = getelementptr inbounds i8, ptr %11, i64 128
  call void @llvm.memset.p0.i64(ptr noundef nonnull align 8 dereferenceable(32) %40, i8 0, i64 32, i1 false)
  call void @_ZN3sbo12small_vectorISt10unique_ptrIN7hipsycl2rt23backend_kernel_launcherESt14default_deleteIS4_EELm8EEC2Ev(ptr noundef nonnull align 8 dereferenceable(96) %12) #7
  invoke void @_ZN7hipsycl2rt15kernel_launcherC2ERKNS_4glue20kernel_launcher_dataEN3sbo12small_vectorISt10unique_ptrINS0_23backend_kernel_launcherESt14default_deleteIS9_EELm8EEE(ptr noundef nonnull align 8 dereferenceable(416) %10, ptr noundef nonnull align 8 dereferenceable(160) %11, ptr noundef nonnull %12)
          to label %41 unwind label %86

41:                                               ; preds = %35
  invoke void @_ZN7hipsycl2rt17requirements_listC2ERKS1_(ptr noundef nonnull align 8 dereferenceable(168) %13, ptr noundef nonnull align 8 dereferenceable(168) %8)
          to label %42 unwind label %88

42:                                               ; preds = %41
  invoke void @_ZN7hipsycl2rt14make_operationINS0_16kernel_operationEJPKcNS0_15kernel_launcherENS0_17requirements_listEEEESt10unique_ptrINS0_9operationESt14default_deleteIS8_EEDpT0_(ptr dead_on_unwind nonnull writable sret(%"class.std::unique_ptr") align 8 %9, ptr noundef nonnull @.str.6, ptr noundef nonnull %10, ptr noundef nonnull %13)
          to label %43 unwind label %90

43:                                               ; preds = %42
  call void @_ZN7hipsycl2rt17requirements_listD2Ev(ptr noundef nonnull align 8 dereferenceable(168) %13) #7
  call void @_ZN7hipsycl2rt15kernel_launcherD2Ev(ptr noundef nonnull align 8 dereferenceable(416) %10) #7
  call void @_ZNSt6vectorISt10unique_ptrIN7hipsycl2rt23backend_kernel_launcherESt14default_deleteIS3_EEN3sbo29small_buffer_vector_allocatorIS6_Lm8ES6_EEED2Ev(ptr noundef nonnull align 8 dereferenceable(96) %12) #7
  call void @_ZN7hipsycl4glue20kernel_launcher_dataD2Ev(ptr noundef nonnull align 8 dereferenceable(160) %11) #7
  call void @llvm.lifetime.end.p0(i64 160, ptr nonnull %11) #7
  call void @llvm.lifetime.start.p0(i64 16, ptr nonnull %14) #7
  call void @_ZNSt10unique_ptrIN7hipsycl2rt9operationESt14default_deleteIS2_EEC2EOS5_(ptr noundef nonnull align 8 dereferenceable(8) %15, ptr noundef nonnull align 8 dereferenceable(8) %9) #7
  invoke void @_ZN7hipsycl2rt11dag_builder17add_command_groupESt10unique_ptrINS0_9operationESt14default_deleteIS3_EERKNS0_17requirements_listERKNS0_15execution_hintsE(ptr dead_on_unwind nonnull writable sret(%"class.std::shared_ptr.53") align 8 %14, ptr noundef nonnull align 8 dereferenceable(368) %7, ptr noundef nonnull %15, ptr noundef nonnull align 8 dereferenceable(168) %8, ptr noundef nonnull align 8 dereferenceable(132) %3)
          to label %44 unwind label %94

44:                                               ; preds = %43
  call void @_ZNSt10unique_ptrIN7hipsycl2rt9operationESt14default_deleteIS2_EED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %15) #7
  %45 = call noundef ptr @_ZNKSt19__shared_ptr_accessIN7hipsycl2rt8dag_nodeELN9__gnu_cxx12_Lock_policyE2ELb0ELb0EEptEv(ptr noundef nonnull align 1 dereferenceable(1) %14) #7
  %46 = invoke noundef nonnull align 8 dereferenceable(132) ptr @_ZN7hipsycl2rt8dag_node19get_execution_hintsEv(ptr noundef nonnull align 8 dereferenceable(384) %45)
          to label %.preheader33 unwind label %.loopexit.split-lp.loopexit.split-lp

.preheader33:                                     ; preds = %44, %56
  %47 = load ptr, ptr @_ZN5boost9unit_test12_GLOBAL__N_113unit_test_logE, align 8, !tbaa !22
  call void @_ZN5boost9unit_test13basic_cstringIKcEC2ImEEPS2_T_(ptr noundef nonnull align 8 dereferenceable(16) %16, ptr noundef nonnull @.str.3, i64 noundef 86)
  invoke void @_ZN5boost9unit_test13basic_cstringIKcEC2Ev(ptr noundef nonnull align 8 dereferenceable(16) %17)
          to label %48 unwind label %.loopexit.split-lp.loopexit

48:                                               ; preds = %.preheader33
  invoke void @_ZN5boost9unit_test15unit_test_log_t14set_checkpointENS0_13basic_cstringIKcEEmS4_(ptr noundef nonnull align 8 dereferenceable(8) %47, ptr noundef nonnull %16, i64 noundef 48, ptr noundef nonnull %17)
          to label %49 unwind label %.loopexit.split-lp.loopexit

49:                                               ; preds = %48
  call void @llvm.lifetime.start.p0(i64 24, ptr nonnull %18) #7
  %50 = invoke noundef zeroext i1 @_ZNK7hipsycl2rt15execution_hints8has_hintINS0_5hints14bind_to_deviceEEEbv(ptr noundef nonnull align 8 dereferenceable(132) %46)
          to label %51 unwind label %96

51:                                               ; preds = %49
  invoke void @_ZN5boost10test_tools16assertion_resultC2Eb(ptr noundef nonnull align 8 dereferenceable(24) %18, i1 noundef zeroext %50)
          to label %52 unwind label %96

52:                                               ; preds = %51
  call void @llvm.lifetime.start.p0(i64 32, ptr nonnull %19) #7
  %53 = call noundef nonnull align 8 dereferenceable(9) ptr @_ZN5boost9unit_test12lazy_ostream8instanceEv()
  call void @llvm.lifetime.start.p0(i64 16, ptr nonnull %20) #7
  call void @_ZN5boost9unit_test13basic_cstringIKcEC2ImEEPS2_T_(ptr noundef nonnull align 8 dereferenceable(16) %20, ptr noundef nonnull @.str.7, i64 noundef 48)
  invoke void @_ZN5boost9unit_testlsINS0_13basic_cstringIKcEEEENS0_17lazy_ostream_implINS0_12lazy_ostreamET_RKS7_EERKS6_S9_(ptr dead_on_unwind nonnull writable sret(%"class.boost::unit_test::lazy_ostream_impl") align 8 %19, ptr noundef nonnull align 8 dereferenceable(9) %53, ptr noundef nonnull align 8 dereferenceable(16) %20)
          to label %54 unwind label %98

54:                                               ; preds = %52
  call void @_ZN5boost9unit_test13basic_cstringIKcEC2ImEEPS2_T_(ptr noundef nonnull align 8 dereferenceable(16) %21, ptr noundef nonnull @.str.3, i64 noundef 86)
  %55 = invoke noundef zeroext i1 (ptr, ptr, ptr, i64, i32, i32, i64, ...) @_ZN5boost10test_tools9tt_detail16report_assertionERKNS0_16assertion_resultERKNS_9unit_test12lazy_ostreamENS5_13basic_cstringIKcEEmNS1_10tool_levelENS1_10check_typeEmz(ptr noundef nonnull align 8 dereferenceable(24) %18, ptr noundef nonnull align 8 dereferenceable(9) %19, ptr noundef nonnull %21, i64 noundef 48, i32 noundef 1, i32 noundef 0, i64 noundef 0)
          to label %56 unwind label %100

56:                                               ; preds = %54
  call void @_ZN5boost9unit_test12lazy_ostreamD2Ev(ptr noundef nonnull align 8 dereferenceable(32) %19) #7
  call void @llvm.lifetime.end.p0(i64 16, ptr nonnull %20) #7
  call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %19) #7
  call void @_ZN5boost10test_tools16assertion_resultD2Ev(ptr noundef nonnull align 8 dereferenceable(24) %18) #7
  call void @llvm.lifetime.end.p0(i64 24, ptr nonnull %18) #7
  %57 = call noundef zeroext i1 @_ZN5boost10test_tools9tt_detail10dummy_condEv()
  br i1 %57, label %.preheader33, label %.preheader, !llvm.loop !23

.preheader:                                       ; preds = %56
  %58 = getelementptr inbounds i8, ptr %25, i64 8
  br label %59

59:                                               ; preds = %.preheader, %73
  %60 = load ptr, ptr @_ZN5boost9unit_test12_GLOBAL__N_113unit_test_logE, align 8, !tbaa !22
  call void @_ZN5boost9unit_test13basic_cstringIKcEC2ImEEPS2_T_(ptr noundef nonnull align 8 dereferenceable(16) %22, ptr noundef nonnull @.str.3, i64 noundef 86)
  invoke void @_ZN5boost9unit_test13basic_cstringIKcEC2Ev(ptr noundef nonnull align 8 dereferenceable(16) %23)
          to label %61 unwind label %.loopexit

61:                                               ; preds = %59
  invoke void @_ZN5boost9unit_test15unit_test_log_t14set_checkpointENS0_13basic_cstringIKcEEmS4_(ptr noundef nonnull align 8 dereferenceable(8) %60, ptr noundef nonnull %22, i64 noundef 50, ptr noundef nonnull %23)
          to label %62 unwind label %.loopexit

62:                                               ; preds = %61
  call void @llvm.lifetime.start.p0(i64 24, ptr nonnull %24) #7
  call void @llvm.lifetime.start.p0(i64 16, ptr nonnull %25) #7
  %63 = call noundef ptr @_ZNK7hipsycl2rt15execution_hints8get_hintINS0_5hints14bind_to_deviceETnNSt9enable_ifIXsr3stdE12is_base_of_vINS3_14execution_hintET_EEiE4typeELi0EEEPKS7_v(ptr noundef nonnull align 8 dereferenceable(132) %46)
  %64 = call { i64, i64 } @_ZNK7hipsycl2rt5hints14bind_to_device13get_device_idEv(ptr noundef nonnull align 4 dereferenceable(20) %63)
  %65 = extractvalue { i64, i64 } %64, 0
  store i64 %65, ptr %25, align 8
  %66 = extractvalue { i64, i64 } %64, 1
  store i64 %66, ptr %58, align 8
  %67 = invoke noundef zeroext i1 @_ZN7hipsycl2rteqERKNS0_9device_idES3_(ptr noundef nonnull align 4 dereferenceable(16) %25, ptr noundef nonnull align 4 dereferenceable(16) %4)
          to label %68 unwind label %104

68:                                               ; preds = %62
  invoke void @_ZN5boost10test_tools16assertion_resultC2Eb(ptr noundef nonnull align 8 dereferenceable(24) %24, i1 noundef zeroext %67)
          to label %69 unwind label %104

69:                                               ; preds = %68
  call void @llvm.lifetime.start.p0(i64 32, ptr nonnull %26) #7
  %70 = call noundef nonnull align 8 dereferenceable(9) ptr @_ZN5boost9unit_test12lazy_ostream8instanceEv()
  call void @llvm.lifetime.start.p0(i64 16, ptr nonnull %27) #7
  call void @_ZN5boost9unit_test13basic_cstringIKcEC2ImEEPS2_T_(ptr noundef nonnull align 8 dereferenceable(16) %27, ptr noundef nonnull @.str.8, i64 noundef 71)
  invoke void @_ZN5boost9unit_testlsINS0_13basic_cstringIKcEEEENS0_17lazy_ostream_implINS0_12lazy_ostreamET_RKS7_EERKS6_S9_(ptr dead_on_unwind nonnull writable sret(%"class.boost::unit_test::lazy_ostream_impl") align 8 %26, ptr noundef nonnull align 8 dereferenceable(9) %70, ptr noundef nonnull align 8 dereferenceable(16) %27)
          to label %71 unwind label %106

71:                                               ; preds = %69
  call void @_ZN5boost9unit_test13basic_cstringIKcEC2ImEEPS2_T_(ptr noundef nonnull align 8 dereferenceable(16) %28, ptr noundef nonnull @.str.3, i64 noundef 86)
  %72 = invoke noundef zeroext i1 (ptr, ptr, ptr, i64, i32, i32, i64, ...) @_ZN5boost10test_tools9tt_detail16report_assertionERKNS0_16assertion_resultERKNS_9unit_test12lazy_ostreamENS5_13basic_cstringIKcEEmNS1_10tool_levelENS1_10check_typeEmz(ptr noundef nonnull align 8 dereferenceable(24) %24, ptr noundef nonnull align 8 dereferenceable(9) %26, ptr noundef nonnull %28, i64 noundef 50, i32 noundef 1, i32 noundef 0, i64 noundef 0)
          to label %73 unwind label %108

73:                                               ; preds = %71
  call void @_ZN5boost9unit_test12lazy_ostreamD2Ev(ptr noundef nonnull align 8 dereferenceable(32) %26) #7
  call void @llvm.lifetime.end.p0(i64 16, ptr nonnull %27) #7
  call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %26) #7
  call void @_ZN5boost10test_tools16assertion_resultD2Ev(ptr noundef nonnull align 8 dereferenceable(24) %24) #7
  call void @llvm.lifetime.end.p0(i64 16, ptr nonnull %25) #7
  call void @llvm.lifetime.end.p0(i64 24, ptr nonnull %24) #7
  %74 = call noundef zeroext i1 @_ZN5boost10test_tools9tt_detail10dummy_condEv()
  br i1 %74, label %59, label %75, !llvm.loop !25

75:                                               ; preds = %73
  %76 = call noundef ptr @_ZNKSt19__shared_ptr_accessIN7hipsycl2rt8dag_nodeELN9__gnu_cxx12_Lock_policyE2ELb0ELb0EEptEv(ptr noundef nonnull align 1 dereferenceable(1) %14) #7
  invoke void @_ZN7hipsycl2rt8dag_node6cancelEv(ptr noundef nonnull align 8 dereferenceable(384) %76)
          to label %77 unwind label %.loopexit.split-lp.loopexit.split-lp

77:                                               ; preds = %75
  call void @_ZNSt12__shared_ptrIN7hipsycl2rt8dag_nodeELN9__gnu_cxx12_Lock_policyE2EED2Ev(ptr noundef nonnull align 8 dereferenceable(16) %14) #7
  call void @llvm.lifetime.end.p0(i64 16, ptr nonnull %14) #7
  call void @_ZNSt10unique_ptrIN7hipsycl2rt9operationESt14default_deleteIS2_EED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %9) #7
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %9) #7
  call void @_ZN7hipsycl2rt17requirements_listD2Ev(ptr noundef nonnull align 8 dereferenceable(168) %8) #7
  call void @llvm.lifetime.end.p0(i64 168, ptr nonnull %8) #7
  call void @_ZN7hipsycl2rt11dag_builderD2Ev(ptr noundef nonnull align 8 dereferenceable(368) %7) #7
  call void @llvm.lifetime.end.p0(i64 368, ptr nonnull %7) #7
  call void @llvm.lifetime.end.p0(i64 16, ptr nonnull %4) #7
  call void @_ZN7hipsycl2rt15execution_hintsD2Ev(ptr noundef nonnull align 8 dereferenceable(132) %3) #7
  call void @llvm.lifetime.end.p0(i64 136, ptr nonnull %3) #7
  call void @_ZN7hipsycl2rt24runtime_keep_alive_tokenD2Ev(ptr noundef nonnull align 8 dereferenceable(16) %2) #7
  call void @llvm.lifetime.end.p0(i64 16, ptr nonnull %2) #7
  ret void

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

80:                                               ; preds = %29
  %81 = landingpad { ptr, i32 }
          cleanup
  call void @llvm.lifetime.end.p0(i64 20, ptr nonnull %6) #7
  br label %116

82:                                               ; preds = %32, %30
  %83 = landingpad { ptr, i32 }
          cleanup
  br label %115

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

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

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

90:                                               ; preds = %42
  %91 = landingpad { ptr, i32 }
          cleanup
  call void @_ZN7hipsycl2rt17requirements_listD2Ev(ptr noundef nonnull align 8 dereferenceable(168) %13) #7
  br label %92

92:                                               ; preds = %90, %88
  %.pn = phi { ptr, i32 } [ %91, %90 ], [ %89, %88 ]
  call void @_ZN7hipsycl2rt15kernel_launcherD2Ev(ptr noundef nonnull align 8 dereferenceable(416) %10) #7
  br label %93

93:                                               ; preds = %92, %86
  %.pn.pn = phi { ptr, i32 } [ %.pn, %92 ], [ %87, %86 ]
  call void @_ZNSt6vectorISt10unique_ptrIN7hipsycl2rt23backend_kernel_launcherESt14default_deleteIS3_EEN3sbo29small_buffer_vector_allocatorIS6_Lm8ES6_EEED2Ev(ptr noundef nonnull align 8 dereferenceable(96) %12) #7
  call void @_ZN7hipsycl4glue20kernel_launcher_dataD2Ev(ptr noundef nonnull align 8 dereferenceable(160) %11) #7
  call void @llvm.lifetime.end.p0(i64 160, ptr nonnull %11) #7
  br label %113

94:                                               ; preds = %43
  %95 = landingpad { ptr, i32 }
          cleanup
  call void @_ZNSt10unique_ptrIN7hipsycl2rt9operationESt14default_deleteIS2_EED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %15) #7
  br label %112

.loopexit:                                        ; preds = %59, %61
  %lpad.loopexit = landingpad { ptr, i32 }
          cleanup
  br label %.loopexit.split-lp

.loopexit.split-lp.loopexit:                      ; preds = %48, %.preheader33
  %lpad.loopexit34 = landingpad { ptr, i32 }
          cleanup
  br label %.loopexit.split-lp

.loopexit.split-lp.loopexit.split-lp:             ; preds = %75, %44
  %lpad.loopexit.split-lp35 = landingpad { ptr, i32 }
          cleanup
  br label %.loopexit.split-lp

96:                                               ; preds = %51, %49
  %97 = landingpad { ptr, i32 }
          cleanup
  br label %103

98:                                               ; preds = %52
  %99 = landingpad { ptr, i32 }
          cleanup
  br label %102

100:                                              ; preds = %54
  %101 = landingpad { ptr, i32 }
          cleanup
  call void @_ZN5boost9unit_test12lazy_ostreamD2Ev(ptr noundef nonnull align 8 dereferenceable(32) %19) #7
  br label %102

102:                                              ; preds = %100, %98
  %.pn20 = phi { ptr, i32 } [ %101, %100 ], [ %99, %98 ]
  call void @llvm.lifetime.end.p0(i64 16, ptr nonnull %20) #7
  call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %19) #7
  call void @_ZN5boost10test_tools16assertion_resultD2Ev(ptr noundef nonnull align 8 dereferenceable(24) %18) #7
  br label %103

103:                                              ; preds = %102, %96
  %.pn20.pn = phi { ptr, i32 } [ %.pn20, %102 ], [ %97, %96 ]
  call void @llvm.lifetime.end.p0(i64 24, ptr nonnull %18) #7
  br label %.loopexit.split-lp

104:                                              ; preds = %68, %62
  %105 = landingpad { ptr, i32 }
          cleanup
  br label %111

106:                                              ; preds = %69
  %107 = landingpad { ptr, i32 }
          cleanup
  br label %110

108:                                              ; preds = %71
  %109 = landingpad { ptr, i32 }
          cleanup
  call void @_ZN5boost9unit_test12lazy_ostreamD2Ev(ptr noundef nonnull align 8 dereferenceable(32) %26) #7
  br label %110

110:                                              ; preds = %108, %106
  %.pn23 = phi { ptr, i32 } [ %109, %108 ], [ %107, %106 ]
  call void @llvm.lifetime.end.p0(i64 16, ptr nonnull %27) #7
  call void @llvm.lifetime.end.p0(i64 32, ptr nonnull %26) #7
  call void @_ZN5boost10test_tools16assertion_resultD2Ev(ptr noundef nonnull align 8 dereferenceable(24) %24) #7
  br label %111

111:                                              ; preds = %110, %104
  %.pn23.pn = phi { ptr, i32 } [ %.pn23, %110 ], [ %105, %104 ]
  call void @llvm.lifetime.end.p0(i64 16, ptr nonnull %25) #7
  call void @llvm.lifetime.end.p0(i64 24, ptr nonnull %24) #7
  br label %.loopexit.split-lp

.loopexit.split-lp:                               ; preds = %.loopexit, %.loopexit.split-lp.loopexit.split-lp, %.loopexit.split-lp.loopexit, %111, %103
  %.pn26 = phi { ptr, i32 } [ %.pn23.pn, %111 ], [ %.pn20.pn, %103 ], [ %lpad.loopexit, %.loopexit ], [ %lpad.loopexit34, %.loopexit.split-lp.loopexit ], [ %lpad.loopexit.split-lp35, %.loopexit.split-lp.loopexit.split-lp ]
  call void @_ZNSt12__shared_ptrIN7hipsycl2rt8dag_nodeELN9__gnu_cxx12_Lock_policyE2EED2Ev(ptr noundef nonnull align 8 dereferenceable(16) %14) #7
  br label %112

112:                                              ; preds = %.loopexit.split-lp, %94
  %.pn26.pn = phi { ptr, i32 } [ %.pn26, %.loopexit.split-lp ], [ %95, %94 ]
  call void @llvm.lifetime.end.p0(i64 16, ptr nonnull %14) #7
  call void @_ZNSt10unique_ptrIN7hipsycl2rt9operationESt14default_deleteIS2_EED2Ev(ptr noundef nonnull align 8 dereferenceable(8) %9) #7
  br label %113

113:                                              ; preds = %112, %93
  %.pn26.pn.pn = phi { ptr, i32 } [ %.pn26.pn, %112 ], [ %.pn.pn, %93 ]
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %9) #7
  call void @_ZN7hipsycl2rt17requirements_listD2Ev(ptr noundef nonnull align 8 dereferenceable(168) %8) #7
  br label %114

114:                                              ; preds = %113, %84
  %.pn26.pn.pn.pn = phi { ptr, i32 } [ %.pn26.pn.pn, %113 ], [ %85, %84 ]
  call void @llvm.lifetime.end.p0(i64 168, ptr nonnull %8) #7
  call void @_ZN7hipsycl2rt11dag_builderD2Ev(ptr noundef nonnull align 8 dereferenceable(368) %7) #7
  br label %115

115:                                              ; preds = %114, %82
  %.pn26.pn.pn.pn.pn = phi { ptr, i32 } [ %.pn26.pn.pn.pn, %114 ], [ %83, %82 ]
  call void @llvm.lifetime.end.p0(i64 368, ptr nonnull %7) #7
  br label %116

116:                                              ; preds = %115, %80, %78
  %.pn26.pn.pn.pn.pn.pn = phi { ptr, i32 } [ %.pn26.pn.pn.pn.pn, %115 ], [ %81, %80 ], [ %79, %78 ]
  call void @llvm.lifetime.end.p0(i64 16, ptr nonnull %4) #7
  call void @_ZN7hipsycl2rt15execution_hintsD2Ev(ptr noundef nonnull align 8 dereferenceable(132) %3) #7
  call void @llvm.lifetime.end.p0(i64 136, ptr nonnull %3) #7
  call void @_ZN7hipsycl2rt24runtime_keep_alive_tokenD2Ev(ptr noundef nonnull align 8 dereferenceable(16) %2) #7
  call void @llvm.lifetime.end.p0(i64 16, ptr nonnull %2) #7
  resume { ptr, i32 } %.pn26.pn.pn.pn.pn.pn
}

declare void @_ZN7hipsycl2rt24runtime_keep_alive_tokenC1Ev(ptr noundef nonnull align 8 dereferenceable(16)) unnamed_addr #2

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

; Function Attrs: mustprogress nounwind uwtable
declare dso_local void @_ZN7hipsycl2rt18backend_descriptorC2ENS0_17hardware_platformENS0_12api_platformE(ptr noundef nonnull align 4 dereferenceable(12), i32 noundef, i32 noundef) unnamed_addr #3 align 2

declare void @_ZN7hipsycl2rt9device_idC1ENS0_18backend_descriptorEi(ptr noundef nonnull align 4 dereferenceable(16), i64, i32, i32 noundef) unnamed_addr #2

; Function Attrs: mustprogress uwtable
declare dso_local void @_ZN7hipsycl2rt15execution_hints8set_hintINS0_5hints14bind_to_deviceETnNSt9enable_ifIXsr3stdE12is_base_of_vINS3_14execution_hintET_EEiE4typeELi0EEEvOS7_(ptr noundef nonnull align 8 dereferenceable(132), ptr noundef nonnull align 4 dereferenceable(20)) local_unnamed_addr #1 align 2

; Function Attrs: mustprogress nounwind uwtable
declare dso_local void @_ZN7hipsycl2rt5hints14bind_to_deviceC2ENS0_9device_idE(ptr noundef nonnull align 4 dereferenceable(20), i64, i64) unnamed_addr #3 align 2

declare noundef ptr @_ZNK7hipsycl2rt24runtime_keep_alive_token3getEv(ptr noundef nonnull align 8 dereferenceable(16)) local_unnamed_addr #2

declare void @_ZN7hipsycl2rt11dag_builderC1EPNS0_7runtimeE(ptr noundef nonnull align 8 dereferenceable(368), ptr noundef) unnamed_addr #2

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

; Function Attrs: mustprogress uwtable
declare dso_local void @_ZN7hipsycl2rt14make_operationINS0_16kernel_operationEJPKcNS0_15kernel_launcherENS0_17requirements_listEEEESt10unique_ptrINS0_9operationESt14default_deleteIS8_EEDpT0_(ptr dead_on_unwind noalias writable sret(%"class.std::unique_ptr") align 8, ptr noundef, ptr noundef, ptr noundef) local_unnamed_addr #1

; Function Attrs: mustprogress nocallback nofree nounwind willreturn memory(argmem: write)
declare void @llvm.memset.p0.i64(ptr nocapture writeonly, i8, i64, i1 immarg) #4

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

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

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

; Function Attrs: mustprogress uwtable
declare dso_local void @_ZN7hipsycl2rt15kernel_launcherC2ERKNS_4glue20kernel_launcher_dataEN3sbo12small_vectorISt10unique_ptrINS0_23backend_kernel_launcherESt14default_deleteIS9_EELm8EEE(ptr noundef nonnull align 8 dereferenceable(416), ptr noundef nonnull align 8 dereferenceable(160), ptr noundef) unnamed_addr #1 align 2

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

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

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

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

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

declare void @_ZN7hipsycl2rt11dag_builder17add_command_groupESt10unique_ptrINS0_9operationESt14default_deleteIS3_EERKNS0_17requirements_listERKNS0_15execution_hintsE(ptr dead_on_unwind writable sret(%"class.std::shared_ptr.53") align 8, ptr noundef nonnull align 8 dereferenceable(368), ptr noundef, ptr noundef nonnull align 8 dereferenceable(168), ptr noundef nonnull align 8 dereferenceable(132)) local_unnamed_addr #2

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

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

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

declare noundef nonnull align 8 dereferenceable(132) ptr @_ZN7hipsycl2rt8dag_node19get_execution_hintsEv(ptr noundef nonnull align 8 dereferenceable(384)) local_unnamed_addr #2

declare void @_ZN5boost9unit_test15unit_test_log_t14set_checkpointENS0_13basic_cstringIKcEEmS4_(ptr noundef nonnull align 8 dereferenceable(8), ptr noundef, i64 noundef, ptr noundef) local_unnamed_addr #2

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

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

declare noundef zeroext i1 @_ZN5boost10test_tools9tt_detail16report_assertionERKNS0_16assertion_resultERKNS_9unit_test12lazy_ostreamENS5_13basic_cstringIKcEEmNS1_10tool_levelENS1_10check_typeEmz(ptr noundef nonnull align 8 dereferenceable(24), ptr noundef nonnull align 8 dereferenceable(9), ptr noundef, i64 noundef, i32 noundef, i32 noundef, i64 noundef, ...) local_unnamed_addr #2

; Function Attrs: mustprogress uwtable
declare dso_local noundef zeroext i1 @_ZNK7hipsycl2rt15execution_hints8has_hintINS0_5hints14bind_to_deviceEEEbv(ptr noundef nonnull align 8 dereferenceable(132)) local_unnamed_addr #1 align 2

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

; Function Attrs: inlinehint mustprogress uwtable
declare dso_local void @_ZN5boost9unit_testlsINS0_13basic_cstringIKcEEEENS0_17lazy_ostream_implINS0_12lazy_ostreamET_RKS7_EERKS6_S9_(ptr dead_on_unwind noalias writable sret(%"class.boost::unit_test::lazy_ostream_impl") align 8, ptr noundef nonnull align 8 dereferenceable(9), ptr noundef nonnull align 8 dereferenceable(16)) local_unnamed_addr #5

; Function Attrs: mustprogress nounwind uwtable
declare dso_local noundef nonnull align 8 dereferenceable(9) ptr @_ZN5boost9unit_test12lazy_ostream8instanceEv() local_unnamed_addr #3 align 2

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

; Function Attrs: inlinehint mustprogress nounwind uwtable
declare dso_local noundef zeroext i1 @_ZN5boost10test_tools9tt_detail10dummy_condEv() local_unnamed_addr #6

; Function Attrs: mustprogress uwtable
declare dso_local noundef zeroext i1 @_ZN7hipsycl2rteqERKNS0_9device_idES3_(ptr noundef nonnull align 4 dereferenceable(16), ptr noundef nonnull align 4 dereferenceable(16)) local_unnamed_addr #1

; Function Attrs: mustprogress nounwind uwtable
declare dso_local noundef ptr @_ZNK7hipsycl2rt15execution_hints8get_hintINS0_5hints14bind_to_deviceETnNSt9enable_ifIXsr3stdE12is_base_of_vINS3_14execution_hintET_EEiE4typeELi0EEEPKS7_v(ptr noundef nonnull align 8 dereferenceable(132)) local_unnamed_addr #3 align 2

; Function Attrs: mustprogress nounwind uwtable
declare dso_local { i64, i64 } @_ZNK7hipsycl2rt5hints14bind_to_device13get_device_idEv(ptr noundef nonnull align 4 dereferenceable(20)) local_unnamed_addr #3 align 2

declare void @_ZN7hipsycl2rt8dag_node6cancelEv(ptr noundef nonnull align 8 dereferenceable(384)) local_unnamed_addr #2

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

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

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

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

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

attributes #0 = { mustprogress nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) }
attributes #1 = { 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 #2 = { "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cmov,+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #3 = { mustprogress 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 = { mustprogress nocallback nofree nounwind willreturn memory(argmem: write) }
attributes #5 = { inlinehint mustprogress uwtable "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cmov,+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #6 = { 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 #7 = { 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, !7, i64 0}
!6 = !{!"_ZTSN7hipsycl4glue20kernel_launcher_dataE", !7, i64 0, !10, i64 8, !11, i64 16, !15, i64 40, !15, i64 64, !17, i64 88, !18, i64 96, !20, i64 128, !7, i64 136, !7, i64 144, !7, i64 152}
!7 = !{!"any pointer", !8, i64 0}
!8 = !{!"omnipotent char", !9, i64 0}
!9 = !{!"Simple C++ TBAA"}
!10 = !{!"_ZTSN7hipsycl2rt11kernel_typeE", !8, i64 0}
!11 = !{!"_ZTSSt6vectorIhSaIhEE", !12, i64 0}
!12 = !{!"_ZTSSt12_Vector_baseIhSaIhEE", !13, i64 0}
!13 = !{!"_ZTSNSt12_Vector_baseIhSaIhEE12_Vector_implE", !14, i64 0}
!14 = !{!"_ZTSNSt12_Vector_baseIhSaIhEE17_Vector_impl_dataE", !7, i64 0, !7, i64 8, !7, i64 16}
!15 = !{!"_ZTSN7hipsycl2rt12static_arrayILi3EEE", !16, i64 0}
!16 = !{!"_ZTSSt5arrayImLm3EE", !8, i64 0}
!17 = !{!"int", !8, i64 0}
!18 = !{!"_ZTSSt8functionIFvPN7hipsycl2rt16kernel_operationERNS0_4sycl14interop_handleEEE", !19, i64 0, !7, i64 24}
!19 = !{!"_ZTSSt14_Function_base", !8, i64 0, !7, i64 16}
!20 = !{!"long long", !8, i64 0}
!21 = !{!6, !10, i64 8}
!22 = !{!7, !7, i64 0}
!23 = distinct !{!23, !24}
!24 = !{!"llvm.loop.mustprogress"}
!25 = distinct !{!25, !24}


; Effective O3 passes: EarlyCSEPass,GlobalOptPass,InferAlignmentPass,InferFunctionAttrsPass,InstCombinePass,LICMPass,LoopSimplifyPass,LoopUnrollPass,MemCpyOptPass,PostOrderFunctionAttrsPass,SROAPass,SimplifyCFGPass