; LLVM IR from AdaptiveCpp/AdaptiveCpp :: src/runtime/pcuda (pcuda_runtime_api)

; 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.std::shared_ptr.150" = type { %"class.std::__shared_ptr.151" }
%"class.std::__shared_ptr.151" = 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.119" }
%"class.std::vector.119" = type { %"struct.std::_Vector_base.120" }
%"struct.std::_Vector_base.120" = 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.156" }
%"class.std::shared_ptr.156" = type { %"class.std::__shared_ptr.157" }
%"class.std::__shared_ptr.157" = 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.sbo::small_vector" = type { %"class.std::vector.159" }
%"class.std::vector.159" = type { %"struct.std::_Vector_base.160" }
%"struct.std::_Vector_base.160" = 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.std::shared_ptr.153" = type { %"class.std::__shared_ptr.154" }
%"class.std::__shared_ptr.154" = type { ptr, %"class.std::__shared_count" }
%"class.hipsycl::rt::result" = type { %"class.std::unique_ptr" }
%"class.std::unique_ptr" = type { %"struct.std::__uniq_ptr_data" }
%"struct.std::__uniq_ptr_data" = type { %"class.std::__uniq_ptr_impl" }
%"class.std::__uniq_ptr_impl" = type { %"class.std::tuple" }
%"class.std::tuple" = type { %"struct.std::_Tuple_impl" }
%"struct.std::_Tuple_impl" = type { %"struct.std::_Head_base.67" }
%"struct.std::_Head_base.67" = type { ptr }

@.str = external hidden unnamed_addr constant [97 x i8], align 1
@.str.63 = external hidden unnamed_addr constant [2 x i8], align 1
@__PRETTY_FUNCTION__.pcudaStreamWaitEvent = external hidden unnamed_addr constant [97 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 noundef nonnull align 8 dereferenceable(104) ptr @_ZN7hipsycl2rt5pcuda17pcuda_application3getEv() #1

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

declare noundef i32 @_ZN7hipsycl2rt5pcuda27get_most_recent_pcuda_errorEv() #1

declare i32 @__gxx_personality_v0(...)

; Function Attrs: mustprogress uwtable
declare hidden noundef ptr @_ZN7hipsycl2rt5pcuda12_GLOBAL__N_122queue_or_default_queueEPNS1_6streamE(ptr noundef) #2

declare noundef zeroext i1 @_ZNK7hipsycl2rt6result10is_successEv(ptr noundef nonnull align 8 dereferenceable(8)) #1

declare void @_ZN7hipsycl2rt5pcuda20register_pcuda_errorERKNS0_6resultE10pcudaError(ptr noundef nonnull align 8 dereferenceable(8), i32 noundef) #1

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

declare noundef nonnull align 8 dereferenceable(40) ptr @_ZN7hipsycl2rt5pcuda17pcuda_application8pcuda_rtEv(ptr noundef nonnull align 8 dereferenceable(104)) #1

; Function Attrs: mustprogress uwtable
declare dso_local noundef ptr @_ZNK7hipsycl2rt5pcuda13pcuda_runtime6get_rtEv(ptr noundef nonnull align 8 dereferenceable(40)) #2 align 2

; Function Attrs: mustprogress nounwind uwtable
declare dso_local noundef nonnull align 8 dereferenceable(72) ptr @_ZN7hipsycl2rt7runtime8backendsEv(ptr noundef nonnull align 8 dereferenceable(608)) #4 align 2

declare noundef ptr @_ZNK7hipsycl2rt15backend_manager3getENS0_10backend_idE(ptr noundef nonnull align 8 dereferenceable(72), i32 noundef) #1

declare noundef i32 @_ZNK7hipsycl2rt9device_id11get_backendEv(ptr noundef nonnull align 4 dereferenceable(16)) #1

; Function Attrs: noreturn nounwind
declare void @__assert_fail(ptr noundef, ptr noundef, i32 noundef, ptr noundef) #5

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

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

declare noundef zeroext i1 @_ZNK7hipsycl2rt5pcuda5event11is_recordedEv(ptr noundef nonnull align 8 dereferenceable(40)) #1

; Function Attrs: mustprogress uwtable
define dso_local i32 @pcudaStreamWaitEvent(ptr noundef %0, ptr noundef %1, i32 noundef %2) #2 personality ptr @__gxx_personality_v0 {
  %4 = alloca i32, align 4
  %5 = alloca ptr, align 8
  %6 = alloca ptr, align 8
  %7 = alloca i32, align 4
  %8 = alloca i32, align 4
  %9 = alloca i32, align 4
  %10 = alloca ptr, align 8
  %11 = alloca %"class.std::shared_ptr.150", align 8
  %12 = alloca %"class.hipsycl::rt::execution_hints", align 8
  %13 = alloca %"class.sbo::small_vector", align 8
  %14 = alloca ptr, align 8
  %15 = alloca ptr, align 8
  %16 = alloca ptr, align 8
  %17 = alloca i32, align 4
  %18 = alloca %"class.std::shared_ptr.153", align 8
  %19 = alloca %"class.hipsycl::rt::device_id", align 4
  %20 = alloca i8, align 1
  %21 = alloca %"class.hipsycl::rt::device_id", align 4
  %22 = alloca %"class.hipsycl::rt::device_id", align 4
  %23 = alloca ptr, align 8
  %24 = alloca %"class.hipsycl::rt::device_id", align 4
  %25 = alloca i64, align 8
  %26 = alloca %"class.hipsycl::rt::device_id", align 4
  %27 = alloca i64, align 8
  %28 = alloca %"class.hipsycl::rt::device_id", align 4
  %29 = alloca %"class.hipsycl::rt::result", align 8
  %30 = alloca %"class.hipsycl::rt::result", align 8
  %31 = alloca %"class.hipsycl::rt::result", align 8
  store ptr %0, ptr %5, align 8, !tbaa !5
  store ptr %1, ptr %6, align 8, !tbaa !5
  store i32 %2, ptr %7, align 4, !tbaa !9
  call void @llvm.lifetime.start.p0(i64 4, ptr %8) #8
  %32 = call noundef i32 @_ZN7hipsycl2rt5pcuda27get_most_recent_pcuda_errorEv()
  store i32 %32, ptr %8, align 4, !tbaa !11
  %33 = load i32, ptr %8, align 4, !tbaa !11
  %34 = icmp ne i32 %33, 0
  br i1 %34, label %35, label %37

35:                                               ; preds = %3
  %36 = load i32, ptr %8, align 4, !tbaa !11
  store i32 %36, ptr %4, align 4
  store i32 1, ptr %9, align 4
  br label %286

37:                                               ; preds = %3
  call void @llvm.lifetime.start.p0(i64 8, ptr %10) #8
  %38 = load ptr, ptr %5, align 8, !tbaa !5
  %39 = call noundef ptr @_ZN7hipsycl2rt5pcuda12_GLOBAL__N_122queue_or_default_queueEPNS1_6streamE(ptr noundef %38)
  store ptr %39, ptr %10, align 8, !tbaa !5
  %40 = load ptr, ptr %10, align 8, !tbaa !5
  %41 = icmp ne ptr %40, null
  br i1 %41, label %43, label %42

42:                                               ; preds = %37
  store i32 33, ptr %4, align 4
  store i32 1, ptr %9, align 4
  br label %285

43:                                               ; preds = %37
  %44 = load ptr, ptr %6, align 8, !tbaa !5
  %45 = icmp ne ptr %44, null
  br i1 %45, label %47, label %46

46:                                               ; preds = %43
  store i32 33, ptr %4, align 4
  store i32 1, ptr %9, align 4
  br label %285

47:                                               ; preds = %43
  %48 = load ptr, ptr %6, align 8, !tbaa !5
  %49 = call noundef zeroext i1 @_ZNK7hipsycl2rt5pcuda5event11is_recordedEv(ptr noundef nonnull align 8 dereferenceable(40) %48)
  br i1 %49, label %51, label %50

50:                                               ; preds = %47
  store i32 0, ptr %4, align 4
  store i32 1, ptr %9, align 4
  br label %285

51:                                               ; preds = %47
  call void @llvm.lifetime.start.p0(i64 16, ptr %11) #8
  call void @llvm.lifetime.start.p0(i64 136, ptr %12) #8
  call void @llvm.memset.p0.i64(ptr align 8 %12, i8 0, i64 136, i1 false)
  call void @_ZN7hipsycl2rt15execution_hintsC2Ev(ptr noundef nonnull align 8 dereferenceable(132) %12) #8
  call void @llvm.lifetime.start.p0(i64 160, ptr %13) #8
  call void @_ZN3sbo12small_vectorISt10shared_ptrIN7hipsycl2rt8dag_nodeEELm8EEC2Ev(ptr noundef nonnull align 8 dereferenceable(160) %13) #8
  call void @llvm.lifetime.start.p0(i64 8, ptr %14) #8
  store ptr null, ptr %14, align 8, !tbaa !13
  call void @llvm.lifetime.start.p0(i64 8, ptr %15) #8
  %52 = invoke noundef nonnull align 8 dereferenceable(104) ptr @_ZN7hipsycl2rt5pcuda17pcuda_application3getEv()
          to label %53 unwind label %99

53:                                               ; preds = %51
  %54 = invoke noundef nonnull align 8 dereferenceable(40) ptr @_ZN7hipsycl2rt5pcuda17pcuda_application8pcuda_rtEv(ptr noundef nonnull align 8 dereferenceable(104) %52)
          to label %55 unwind label %99

55:                                               ; preds = %53
  %56 = invoke noundef ptr @_ZNK7hipsycl2rt5pcuda13pcuda_runtime6get_rtEv(ptr noundef nonnull align 8 dereferenceable(40) %54)
          to label %57 unwind label %99

57:                                               ; preds = %55
  store ptr %56, ptr %15, align 8, !tbaa !5
  invoke void @_ZSt11make_sharedIN7hipsycl2rt8dag_nodeEJNS1_15execution_hintsEN3sbo12small_vectorISt10shared_ptrIS2_ELm8EEEDnPNS1_7runtimeEEES6_INSt9enable_ifIXntsr8is_arrayIT_EE5valueESC_E4typeEEDpOT0_(ptr dead_on_unwind writable sret(%"class.std::shared_ptr.150") align 8 %11, ptr noundef nonnull align 8 dereferenceable(132) %12, ptr noundef nonnull align 8 dereferenceable(160) %13, ptr noundef nonnull align 8 dereferenceable(8) %14, ptr noundef nonnull align 8 dereferenceable(8) %15)
          to label %58 unwind label %99

58:                                               ; preds = %57
  call void @llvm.lifetime.end.p0(i64 8, ptr %15) #8
  call void @llvm.lifetime.end.p0(i64 8, ptr %14) #8
  call void @_ZNSt6vectorISt10shared_ptrIN7hipsycl2rt8dag_nodeEEN3sbo29small_buffer_vector_allocatorIS4_Lm8ES4_EEED2Ev(ptr noundef nonnull align 8 dereferenceable(160) %13) #8
  call void @llvm.lifetime.end.p0(i64 160, ptr %13) #8
  call void @_ZN7hipsycl2rt15execution_hintsD2Ev(ptr noundef nonnull align 8 dereferenceable(132) %12) #8
  call void @llvm.lifetime.end.p0(i64 136, ptr %12) #8
  %59 = call noundef ptr @_ZNKSt19__shared_ptr_accessIN7hipsycl2rt8dag_nodeELN9__gnu_cxx12_Lock_policyE2ELb0ELb0EEptEv(ptr noundef nonnull align 1 dereferenceable(1) %11) #8
  %60 = load ptr, ptr %6, align 8, !tbaa !5
  invoke void @_ZNK7hipsycl2rt5pcuda5event20get_event_shared_ptrEv(ptr dead_on_unwind writable sret(%"class.std::shared_ptr.153") align 8 %18, ptr noundef nonnull align 8 dereferenceable(40) %60)
          to label %61 unwind label %103

61:                                               ; preds = %58
  invoke void @_ZN7hipsycl2rt8dag_node14mark_submittedESt10shared_ptrINS0_14dag_node_eventEE(ptr noundef nonnull align 8 dereferenceable(384) %59, ptr noundef %18)
          to label %62 unwind label %107

62:                                               ; preds = %61
  call void @_ZNSt12__shared_ptrIN7hipsycl2rt14dag_node_eventELN9__gnu_cxx12_Lock_policyE2EED2Ev(ptr noundef nonnull align 8 dereferenceable(16) %18) #8
  %63 = call noundef ptr @_ZNKSt19__shared_ptr_accessIN7hipsycl2rt8dag_nodeELN9__gnu_cxx12_Lock_policyE2ELb0ELb0EEptEv(ptr noundef nonnull align 1 dereferenceable(1) %11) #8
  %64 = load ptr, ptr %6, align 8, !tbaa !5
  %65 = invoke { i64, i64 } @_ZNK7hipsycl2rt5pcuda5event10get_deviceEv(ptr noundef nonnull align 8 dereferenceable(40) %64)
          to label %66 unwind label %103

66:                                               ; preds = %62
  %67 = getelementptr inbounds { i64, i64 }, ptr %19, i32 0, i32 0
  %68 = extractvalue { i64, i64 } %65, 0
  store i64 %68, ptr %67, align 4
  %69 = getelementptr inbounds { i64, i64 }, ptr %19, i32 0, i32 1
  %70 = extractvalue { i64, i64 } %65, 1
  store i64 %70, ptr %69, align 4
  %71 = getelementptr inbounds { i64, i64 }, ptr %19, i32 0, i32 0
  %72 = load i64, ptr %71, align 4
  %73 = getelementptr inbounds { i64, i64 }, ptr %19, i32 0, i32 1
  %74 = load i64, ptr %73, align 4
  invoke void @_ZN7hipsycl2rt8dag_node16assign_to_deviceENS0_9device_idE(ptr noundef nonnull align 8 dereferenceable(384) %63, i64 %72, i64 %74)
          to label %75 unwind label %103

75:                                               ; preds = %66
  call void @llvm.lifetime.start.p0(i64 1, ptr %20) #8
  store i8 1, ptr %20, align 1, !tbaa !15
  call void @llvm.lifetime.start.p0(i64 16, ptr %21) #8
  %76 = load ptr, ptr %6, align 8, !tbaa !5
  %77 = invoke { i64, i64 } @_ZNK7hipsycl2rt5pcuda5event10get_deviceEv(ptr noundef nonnull align 8 dereferenceable(40) %76)
          to label %78 unwind label %111

78:                                               ; preds = %75
  %79 = getelementptr inbounds { i64, i64 }, ptr %21, i32 0, i32 0
  %80 = extractvalue { i64, i64 } %77, 0
  store i64 %80, ptr %79, align 4
  %81 = getelementptr inbounds { i64, i64 }, ptr %21, i32 0, i32 1
  %82 = extractvalue { i64, i64 } %77, 1
  store i64 %82, ptr %81, align 4
  %83 = invoke noundef i32 @_ZNK7hipsycl2rt9device_id11get_backendEv(ptr noundef nonnull align 4 dereferenceable(16) %21)
          to label %84 unwind label %111

84:                                               ; preds = %78
  call void @llvm.lifetime.start.p0(i64 16, ptr %22) #8
  %85 = load ptr, ptr %10, align 8, !tbaa !5
  %86 = load ptr, ptr %85, align 8, !tbaa !17
  %87 = getelementptr inbounds ptr, ptr %86, i64 9
  %88 = load ptr, ptr %87, align 8
  %89 = invoke { i64, i64 } %88(ptr noundef nonnull align 8 dereferenceable(8) %85)
          to label %90 unwind label %115

90:                                               ; preds = %84
  %91 = getelementptr inbounds { i64, i64 }, ptr %22, i32 0, i32 0
  %92 = extractvalue { i64, i64 } %89, 0
  store i64 %92, ptr %91, align 4
  %93 = getelementptr inbounds { i64, i64 }, ptr %22, i32 0, i32 1
  %94 = extractvalue { i64, i64 } %89, 1
  store i64 %94, ptr %93, align 4
  %95 = invoke noundef i32 @_ZNK7hipsycl2rt9device_id11get_backendEv(ptr noundef nonnull align 4 dereferenceable(16) %22)
          to label %96 unwind label %115

96:                                               ; preds = %90
  %97 = icmp ne i32 %83, %95
  call void @llvm.lifetime.end.p0(i64 16, ptr %22) #8
  call void @llvm.lifetime.end.p0(i64 16, ptr %21) #8
  br i1 %97, label %98, label %120

98:                                               ; preds = %96
  store i8 0, ptr %20, align 1, !tbaa !15
  br label %233

99:                                               ; preds = %57, %55, %53, %51
  %100 = landingpad { ptr, i32 }
          cleanup
  %101 = extractvalue { ptr, i32 } %100, 0
  store ptr %101, ptr %16, align 8
  %102 = extractvalue { ptr, i32 } %100, 1
  store i32 %102, ptr %17, align 4
  call void @llvm.lifetime.end.p0(i64 8, ptr %15) #8
  call void @llvm.lifetime.end.p0(i64 8, ptr %14) #8
  call void @_ZNSt6vectorISt10shared_ptrIN7hipsycl2rt8dag_nodeEEN3sbo29small_buffer_vector_allocatorIS4_Lm8ES4_EEED2Ev(ptr noundef nonnull align 8 dereferenceable(160) %13) #8
  call void @llvm.lifetime.end.p0(i64 160, ptr %13) #8
  call void @_ZN7hipsycl2rt15execution_hintsD2Ev(ptr noundef nonnull align 8 dereferenceable(132) %12) #8
  call void @llvm.lifetime.end.p0(i64 136, ptr %12) #8
  br label %284

103:                                              ; preds = %66, %62, %58
  %104 = landingpad { ptr, i32 }
          cleanup
  %105 = extractvalue { ptr, i32 } %104, 0
  store ptr %105, ptr %16, align 8
  %106 = extractvalue { ptr, i32 } %104, 1
  store i32 %106, ptr %17, align 4
  br label %283

107:                                              ; preds = %61
  %108 = landingpad { ptr, i32 }
          cleanup
  %109 = extractvalue { ptr, i32 } %108, 0
  store ptr %109, ptr %16, align 8
  %110 = extractvalue { ptr, i32 } %108, 1
  store i32 %110, ptr %17, align 4
  call void @_ZNSt12__shared_ptrIN7hipsycl2rt14dag_node_eventELN9__gnu_cxx12_Lock_policyE2EED2Ev(ptr noundef nonnull align 8 dereferenceable(16) %18) #8
  br label %283

111:                                              ; preds = %78, %75
  %112 = landingpad { ptr, i32 }
          cleanup
  %113 = extractvalue { ptr, i32 } %112, 0
  store ptr %113, ptr %16, align 8
  %114 = extractvalue { ptr, i32 } %112, 1
  store i32 %114, ptr %17, align 4
  br label %119

115:                                              ; preds = %90, %84
  %116 = landingpad { ptr, i32 }
          cleanup
  %117 = extractvalue { ptr, i32 } %116, 0
  store ptr %117, ptr %16, align 8
  %118 = extractvalue { ptr, i32 } %116, 1
  store i32 %118, ptr %17, align 4
  call void @llvm.lifetime.end.p0(i64 16, ptr %22) #8
  br label %119

119:                                              ; preds = %115, %111
  call void @llvm.lifetime.end.p0(i64 16, ptr %21) #8
  br label %282

120:                                              ; preds = %96
  call void @llvm.lifetime.start.p0(i64 8, ptr %23) #8
  %121 = invoke noundef nonnull align 8 dereferenceable(104) ptr @_ZN7hipsycl2rt5pcuda17pcuda_application3getEv()
          to label %122 unwind label %144

122:                                              ; preds = %120
  %123 = invoke noundef nonnull align 8 dereferenceable(40) ptr @_ZN7hipsycl2rt5pcuda17pcuda_application8pcuda_rtEv(ptr noundef nonnull align 8 dereferenceable(104) %121)
          to label %124 unwind label %144

124:                                              ; preds = %122
  %125 = invoke noundef ptr @_ZNK7hipsycl2rt5pcuda13pcuda_runtime6get_rtEv(ptr noundef nonnull align 8 dereferenceable(40) %123)
          to label %126 unwind label %144

126:                                              ; preds = %124
  %127 = invoke noundef nonnull align 8 dereferenceable(72) ptr @_ZN7hipsycl2rt7runtime8backendsEv(ptr noundef nonnull align 8 dereferenceable(608) %125)
          to label %128 unwind label %144

128:                                              ; preds = %126
  call void @llvm.lifetime.start.p0(i64 16, ptr %24) #8
  %129 = load ptr, ptr %6, align 8, !tbaa !5
  %130 = invoke { i64, i64 } @_ZNK7hipsycl2rt5pcuda5event10get_deviceEv(ptr noundef nonnull align 8 dereferenceable(40) %129)
          to label %131 unwind label %148

131:                                              ; preds = %128
  %132 = getelementptr inbounds { i64, i64 }, ptr %24, i32 0, i32 0
  %133 = extractvalue { i64, i64 } %130, 0
  store i64 %133, ptr %132, align 4
  %134 = getelementptr inbounds { i64, i64 }, ptr %24, i32 0, i32 1
  %135 = extractvalue { i64, i64 } %130, 1
  store i64 %135, ptr %134, align 4
  %136 = invoke noundef i32 @_ZNK7hipsycl2rt9device_id11get_backendEv(ptr noundef nonnull align 4 dereferenceable(16) %24)
          to label %137 unwind label %148

137:                                              ; preds = %131
  %138 = invoke noundef ptr @_ZNK7hipsycl2rt15backend_manager3getENS0_10backend_idE(ptr noundef nonnull align 8 dereferenceable(72) %127, i32 noundef %136)
          to label %139 unwind label %148

139:                                              ; preds = %137
  call void @llvm.lifetime.end.p0(i64 16, ptr %24) #8
  store ptr %138, ptr %23, align 8, !tbaa !5
  %140 = load ptr, ptr %23, align 8, !tbaa !5
  %141 = icmp ne ptr %140, null
  br i1 %141, label %142, label %143

142:                                              ; preds = %139
  br label %153

143:                                              ; preds = %139
  call void @__assert_fail(ptr noundef @.str.63, ptr noundef @.str, i32 noundef 849, ptr noundef @__PRETTY_FUNCTION__.pcudaStreamWaitEvent) #9
  unreachable

144:                                              ; preds = %126, %124, %122, %120
  %145 = landingpad { ptr, i32 }
          cleanup
  %146 = extractvalue { ptr, i32 } %145, 0
  store ptr %146, ptr %16, align 8
  %147 = extractvalue { ptr, i32 } %145, 1
  store i32 %147, ptr %17, align 4
  br label %232

148:                                              ; preds = %137, %131, %128
  %149 = landingpad { ptr, i32 }
          cleanup
  %150 = extractvalue { ptr, i32 } %149, 0
  store ptr %150, ptr %16, align 8
  %151 = extractvalue { ptr, i32 } %149, 1
  store i32 %151, ptr %17, align 4
  call void @llvm.lifetime.end.p0(i64 16, ptr %24) #8
  br label %232

152:                                              ; No predecessors!
  br label %153

153:                                              ; preds = %152, %142
  call void @llvm.lifetime.start.p0(i64 8, ptr %25) #8
  %154 = load ptr, ptr %23, align 8, !tbaa !5
  %155 = load ptr, ptr %154, align 8, !tbaa !17
  %156 = getelementptr inbounds ptr, ptr %155, i64 3
  %157 = load ptr, ptr %156, align 8
  %158 = invoke noundef ptr %157(ptr noundef nonnull align 8 dereferenceable(8) %154)
          to label %159 unwind label %213

159:                                              ; preds = %153
  call void @llvm.lifetime.start.p0(i64 16, ptr %26) #8
  %160 = load ptr, ptr %6, align 8, !tbaa !5
  %161 = invoke { i64, i64 } @_ZNK7hipsycl2rt5pcuda5event10get_deviceEv(ptr noundef nonnull align 8 dereferenceable(40) %160)
          to label %162 unwind label %217

162:                                              ; preds = %159
  %163 = getelementptr inbounds { i64, i64 }, ptr %26, i32 0, i32 0
  %164 = extractvalue { i64, i64 } %161, 0
  store i64 %164, ptr %163, align 4
  %165 = getelementptr inbounds { i64, i64 }, ptr %26, i32 0, i32 1
  %166 = extractvalue { i64, i64 } %161, 1
  store i64 %166, ptr %165, align 4
  %167 = invoke noundef i32 @_ZNK7hipsycl2rt9device_id6get_idEv(ptr noundef nonnull align 4 dereferenceable(16) %26)
          to label %168 unwind label %217

168:                                              ; preds = %162
  %169 = sext i32 %167 to i64
  %170 = load ptr, ptr %158, align 8, !tbaa !17
  %171 = getelementptr inbounds ptr, ptr %170, i64 2
  %172 = load ptr, ptr %171, align 8
  %173 = invoke noundef ptr %172(ptr noundef nonnull align 8 dereferenceable(8) %158, i64 noundef %169)
          to label %174 unwind label %217

174:                                              ; preds = %168
  %175 = load ptr, ptr %173, align 8, !tbaa !17
  %176 = getelementptr inbounds ptr, ptr %175, i64 12
  %177 = load ptr, ptr %176, align 8
  %178 = invoke noundef i64 %177(ptr noundef nonnull align 8 dereferenceable(8) %173)
          to label %179 unwind label %217

179:                                              ; preds = %174
  call void @llvm.lifetime.end.p0(i64 16, ptr %26) #8
  store i64 %178, ptr %25, align 8, !tbaa !19
  call void @llvm.lifetime.start.p0(i64 8, ptr %27) #8
  %180 = load ptr, ptr %23, align 8, !tbaa !5
  %181 = load ptr, ptr %180, align 8, !tbaa !17
  %182 = getelementptr inbounds ptr, ptr %181, i64 3
  %183 = load ptr, ptr %182, align 8
  %184 = invoke noundef ptr %183(ptr noundef nonnull align 8 dereferenceable(8) %180)
          to label %185 unwind label %221

185:                                              ; preds = %179
  call void @llvm.lifetime.start.p0(i64 16, ptr %28) #8
  %186 = load ptr, ptr %10, align 8, !tbaa !5
  %187 = load ptr, ptr %186, align 8, !tbaa !17
  %188 = getelementptr inbounds ptr, ptr %187, i64 9
  %189 = load ptr, ptr %188, align 8
  %190 = invoke { i64, i64 } %189(ptr noundef nonnull align 8 dereferenceable(8) %186)
          to label %191 unwind label %225

191:                                              ; preds = %185
  %192 = getelementptr inbounds { i64, i64 }, ptr %28, i32 0, i32 0
  %193 = extractvalue { i64, i64 } %190, 0
  store i64 %193, ptr %192, align 4
  %194 = getelementptr inbounds { i64, i64 }, ptr %28, i32 0, i32 1
  %195 = extractvalue { i64, i64 } %190, 1
  store i64 %195, ptr %194, align 4
  %196 = invoke noundef i32 @_ZNK7hipsycl2rt9device_id6get_idEv(ptr noundef nonnull align 4 dereferenceable(16) %28)
          to label %197 unwind label %225

197:                                              ; preds = %191
  %198 = sext i32 %196 to i64
  %199 = load ptr, ptr %184, align 8, !tbaa !17
  %200 = getelementptr inbounds ptr, ptr %199, i64 2
  %201 = load ptr, ptr %200, align 8
  %202 = invoke noundef ptr %201(ptr noundef nonnull align 8 dereferenceable(8) %184, i64 noundef %198)
          to label %203 unwind label %225

203:                                              ; preds = %197
  %204 = load ptr, ptr %202, align 8, !tbaa !17
  %205 = getelementptr inbounds ptr, ptr %204, i64 12
  %206 = load ptr, ptr %205, align 8
  %207 = invoke noundef i64 %206(ptr noundef nonnull align 8 dereferenceable(8) %202)
          to label %208 unwind label %225

208:                                              ; preds = %203
  call void @llvm.lifetime.end.p0(i64 16, ptr %28) #8
  store i64 %207, ptr %27, align 8, !tbaa !19
  %209 = load i64, ptr %25, align 8, !tbaa !19
  %210 = load i64, ptr %27, align 8, !tbaa !19
  %211 = icmp ne i64 %209, %210
  br i1 %211, label %212, label %229

212:                                              ; preds = %208
  store i8 0, ptr %20, align 1, !tbaa !15
  br label %229

213:                                              ; preds = %153
  %214 = landingpad { ptr, i32 }
          cleanup
  %215 = extractvalue { ptr, i32 } %214, 0
  store ptr %215, ptr %16, align 8
  %216 = extractvalue { ptr, i32 } %214, 1
  store i32 %216, ptr %17, align 4
  br label %231

217:                                              ; preds = %174, %168, %162, %159
  %218 = landingpad { ptr, i32 }
          cleanup
  %219 = extractvalue { ptr, i32 } %218, 0
  store ptr %219, ptr %16, align 8
  %220 = extractvalue { ptr, i32 } %218, 1
  store i32 %220, ptr %17, align 4
  call void @llvm.lifetime.end.p0(i64 16, ptr %26) #8
  br label %231

221:                                              ; preds = %179
  %222 = landingpad { ptr, i32 }
          cleanup
  %223 = extractvalue { ptr, i32 } %222, 0
  store ptr %223, ptr %16, align 8
  %224 = extractvalue { ptr, i32 } %222, 1
  store i32 %224, ptr %17, align 4
  br label %230

225:                                              ; preds = %203, %197, %191, %185
  %226 = landingpad { ptr, i32 }
          cleanup
  %227 = extractvalue { ptr, i32 } %226, 0
  store ptr %227, ptr %16, align 8
  %228 = extractvalue { ptr, i32 } %226, 1
  store i32 %228, ptr %17, align 4
  call void @llvm.lifetime.end.p0(i64 16, ptr %28) #8
  br label %230

229:                                              ; preds = %212, %208
  call void @llvm.lifetime.end.p0(i64 8, ptr %27) #8
  call void @llvm.lifetime.end.p0(i64 8, ptr %25) #8
  call void @llvm.lifetime.end.p0(i64 8, ptr %23) #8
  br label %233

230:                                              ; preds = %225, %221
  call void @llvm.lifetime.end.p0(i64 8, ptr %27) #8
  br label %231

231:                                              ; preds = %230, %217, %213
  call void @llvm.lifetime.end.p0(i64 8, ptr %25) #8
  br label %232

232:                                              ; preds = %231, %148, %144
  call void @llvm.lifetime.end.p0(i64 8, ptr %23) #8
  br label %282

233:                                              ; preds = %229, %98
  call void @llvm.lifetime.start.p0(i64 8, ptr %29) #8
  call void @_ZN7hipsycl2rt6resultC2Ev(ptr noundef nonnull align 8 dereferenceable(8) %29) #8
  %234 = load i8, ptr %20, align 1, !tbaa !15, !range !21, !noundef !22
  %235 = trunc i8 %234 to i1
  br i1 %235, label %236, label %253

236:                                              ; preds = %233
  call void @llvm.lifetime.start.p0(i64 8, ptr %30) #8
  %237 = load ptr, ptr %10, align 8, !tbaa !5
  %238 = load ptr, ptr %237, align 8, !tbaa !17
  %239 = getelementptr inbounds ptr, ptr %238, i64 6
  %240 = load ptr, ptr %239, align 8
  invoke void %240(ptr dead_on_unwind writable sret(%"class.hipsycl::rt::result") align 8 %30, ptr noundef nonnull align 8 dereferenceable(8) %237, ptr noundef nonnull align 8 dereferenceable(16) %11)
          to label %241 unwind label %244

241:                                              ; preds = %236
  %242 = invoke noundef nonnull align 8 dereferenceable(8) ptr @_ZN7hipsycl2rt6resultaSEOS1_(ptr noundef nonnull align 8 dereferenceable(8) %29, ptr noundef nonnull align 8 dereferenceable(8) %30)
          to label %243 unwind label %248

243:                                              ; preds = %241
  call void @_ZN7hipsycl2rt6resultD2Ev(ptr noundef nonnull align 8 dereferenceable(8) %30) #8
  call void @llvm.lifetime.end.p0(i64 8, ptr %30) #8
  br label %270

244:                                              ; preds = %236
  %245 = landingpad { ptr, i32 }
          cleanup
  %246 = extractvalue { ptr, i32 } %245, 0
  store ptr %246, ptr %16, align 8
  %247 = extractvalue { ptr, i32 } %245, 1
  store i32 %247, ptr %17, align 4
  br label %252

248:                                              ; preds = %241
  %249 = landingpad { ptr, i32 }
          cleanup
  %250 = extractvalue { ptr, i32 } %249, 0
  store ptr %250, ptr %16, align 8
  %251 = extractvalue { ptr, i32 } %249, 1
  store i32 %251, ptr %17, align 4
  call void @_ZN7hipsycl2rt6resultD2Ev(ptr noundef nonnull align 8 dereferenceable(8) %30) #8
  br label %252

252:                                              ; preds = %248, %244
  call void @llvm.lifetime.end.p0(i64 8, ptr %30) #8
  br label %281

253:                                              ; preds = %233
  call void @llvm.lifetime.start.p0(i64 8, ptr %31) #8
  %254 = load ptr, ptr %10, align 8, !tbaa !5
  %255 = load ptr, ptr %254, align 8, !tbaa !17
  %256 = getelementptr inbounds ptr, ptr %255, i64 7
  %257 = load ptr, ptr %256, align 8
  invoke void %257(ptr dead_on_unwind writable sret(%"class.hipsycl::rt::result") align 8 %31, ptr noundef nonnull align 8 dereferenceable(8) %254, ptr noundef nonnull align 8 dereferenceable(16) %11)
          to label %258 unwind label %261

258:                                              ; preds = %253
  %259 = invoke noundef nonnull align 8 dereferenceable(8) ptr @_ZN7hipsycl2rt6resultaSEOS1_(ptr noundef nonnull align 8 dereferenceable(8) %29, ptr noundef nonnull align 8 dereferenceable(8) %31)
          to label %260 unwind label %265

260:                                              ; preds = %258
  call void @_ZN7hipsycl2rt6resultD2Ev(ptr noundef nonnull align 8 dereferenceable(8) %31) #8
  call void @llvm.lifetime.end.p0(i64 8, ptr %31) #8
  br label %270

261:                                              ; preds = %253
  %262 = landingpad { ptr, i32 }
          cleanup
  %263 = extractvalue { ptr, i32 } %262, 0
  store ptr %263, ptr %16, align 8
  %264 = extractvalue { ptr, i32 } %262, 1
  store i32 %264, ptr %17, align 4
  br label %269

265:                                              ; preds = %258
  %266 = landingpad { ptr, i32 }
          cleanup
  %267 = extractvalue { ptr, i32 } %266, 0
  store ptr %267, ptr %16, align 8
  %268 = extractvalue { ptr, i32 } %266, 1
  store i32 %268, ptr %17, align 4
  call void @_ZN7hipsycl2rt6resultD2Ev(ptr noundef nonnull align 8 dereferenceable(8) %31) #8
  br label %269

269:                                              ; preds = %265, %261
  call void @llvm.lifetime.end.p0(i64 8, ptr %31) #8
  br label %281

270:                                              ; preds = %260, %243
  %271 = invoke noundef zeroext i1 @_ZNK7hipsycl2rt6result10is_successEv(ptr noundef nonnull align 8 dereferenceable(8) %29)
          to label %272 unwind label %275

272:                                              ; preds = %270
  br i1 %271, label %279, label %273

273:                                              ; preds = %272
  invoke void @_ZN7hipsycl2rt5pcuda20register_pcuda_errorERKNS0_6resultE10pcudaError(ptr noundef nonnull align 8 dereferenceable(8) %29, i32 noundef 30)
          to label %274 unwind label %275

274:                                              ; preds = %273
  store i32 30, ptr %4, align 4
  store i32 1, ptr %9, align 4
  br label %280

275:                                              ; preds = %273, %270
  %276 = landingpad { ptr, i32 }
          cleanup
  %277 = extractvalue { ptr, i32 } %276, 0
  store ptr %277, ptr %16, align 8
  %278 = extractvalue { ptr, i32 } %276, 1
  store i32 %278, ptr %17, align 4
  br label %281

279:                                              ; preds = %272
  store i32 0, ptr %4, align 4
  store i32 1, ptr %9, align 4
  br label %280

280:                                              ; preds = %279, %274
  call void @_ZN7hipsycl2rt6resultD2Ev(ptr noundef nonnull align 8 dereferenceable(8) %29) #8
  call void @llvm.lifetime.end.p0(i64 8, ptr %29) #8
  call void @llvm.lifetime.end.p0(i64 1, ptr %20) #8
  call void @_ZNSt12__shared_ptrIN7hipsycl2rt8dag_nodeELN9__gnu_cxx12_Lock_policyE2EED2Ev(ptr noundef nonnull align 8 dereferenceable(16) %11) #8
  call void @llvm.lifetime.end.p0(i64 16, ptr %11) #8
  br label %285

281:                                              ; preds = %275, %269, %252
  call void @_ZN7hipsycl2rt6resultD2Ev(ptr noundef nonnull align 8 dereferenceable(8) %29) #8
  call void @llvm.lifetime.end.p0(i64 8, ptr %29) #8
  br label %282

282:                                              ; preds = %281, %232, %119
  call void @llvm.lifetime.end.p0(i64 1, ptr %20) #8
  br label %283

283:                                              ; preds = %282, %107, %103
  call void @_ZNSt12__shared_ptrIN7hipsycl2rt8dag_nodeELN9__gnu_cxx12_Lock_policyE2EED2Ev(ptr noundef nonnull align 8 dereferenceable(16) %11) #8
  br label %284

284:                                              ; preds = %283, %99
  call void @llvm.lifetime.end.p0(i64 16, ptr %11) #8
  call void @llvm.lifetime.end.p0(i64 8, ptr %10) #8
  call void @llvm.lifetime.end.p0(i64 4, ptr %8) #8
  br label %288

285:                                              ; preds = %280, %50, %46, %42
  call void @llvm.lifetime.end.p0(i64 8, ptr %10) #8
  br label %286

286:                                              ; preds = %285, %35
  call void @llvm.lifetime.end.p0(i64 4, ptr %8) #8
  %287 = load i32, ptr %4, align 4
  ret i32 %287

288:                                              ; preds = %284
  %289 = load ptr, ptr %16, align 8
  %290 = load i32, ptr %17, align 4
  %291 = insertvalue { ptr, i32 } poison, ptr %289, 0
  %292 = insertvalue { ptr, i32 } %291, i32 %290, 1
  resume { ptr, i32 } %292
}

; Function Attrs: inlinehint mustprogress uwtable
declare dso_local void @_ZSt11make_sharedIN7hipsycl2rt8dag_nodeEJNS1_15execution_hintsEN3sbo12small_vectorISt10shared_ptrIS2_ELm8EEEDnPNS1_7runtimeEEES6_INSt9enable_ifIXntsr8is_arrayIT_EE5valueESC_E4typeEEDpOT0_(ptr dead_on_unwind noalias writable sret(%"class.std::shared_ptr.150") align 8, ptr noundef nonnull align 8 dereferenceable(132), ptr noundef nonnull align 8 dereferenceable(160), ptr noundef nonnull align 8 dereferenceable(8), ptr noundef nonnull align 8 dereferenceable(8)) #7

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

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

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

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

declare void @_ZN7hipsycl2rt8dag_node14mark_submittedESt10shared_ptrINS0_14dag_node_eventEE(ptr noundef nonnull align 8 dereferenceable(384), ptr noundef) #1

; Function Attrs: mustprogress nounwind uwtable
declare dso_local void @_ZNK7hipsycl2rt5pcuda5event20get_event_shared_ptrEv(ptr dead_on_unwind noalias writable sret(%"class.std::shared_ptr.153") align 8, ptr noundef nonnull align 8 dereferenceable(40)) #4 align 2

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

declare void @_ZN7hipsycl2rt8dag_node16assign_to_deviceENS0_9device_idE(ptr noundef nonnull align 8 dereferenceable(384), i64, i64) #1

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

declare noundef i32 @_ZNK7hipsycl2rt9device_id6get_idEv(ptr noundef nonnull align 4 dereferenceable(16)) #1

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

declare noundef nonnull align 8 dereferenceable(8) ptr @_ZN7hipsycl2rt6resultaSEOS1_(ptr noundef nonnull align 8 dereferenceable(8), ptr noundef nonnull align 8 dereferenceable(8)) #1

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

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

!0 = !{i32 1, !"wchar_size", i32 4}
!1 = !{i32 8, !"PIC Level", i32 2}
!2 = !{i32 7, !"PIE Level", i32 2}
!3 = !{i32 7, !"uwtable", i32 2}
!4 = !{!"clang version 19.1.0 (git@github.com:llvm/llvm-project.git a4bf6cd7cfb1a1421ba92bca9d017b49936c55e4)"}
!5 = !{!6, !6, i64 0}
!6 = !{!"any pointer", !7, i64 0}
!7 = !{!"omnipotent char", !8, i64 0}
!8 = !{!"Simple C++ TBAA"}
!9 = !{!10, !10, i64 0}
!10 = !{!"int", !7, i64 0}
!11 = !{!12, !12, i64 0}
!12 = !{!"_ZTS10pcudaError", !7, i64 0}
!13 = !{!14, !14, i64 0}
!14 = !{!"std::nullptr_t", !7, i64 0}
!15 = !{!16, !16, i64 0}
!16 = !{!"bool", !7, i64 0}
!17 = !{!18, !18, i64 0}
!18 = !{!"vtable pointer", !8, i64 0}
!19 = !{!20, !20, i64 0}
!20 = !{!"long", !7, i64 0}
!21 = !{i8 0, i8 2}
!22 = !{}


; 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.std::shared_ptr.150" = type { %"class.std::__shared_ptr.151" }
%"class.std::__shared_ptr.151" = 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.119" }
%"class.std::vector.119" = type { %"struct.std::_Vector_base.120" }
%"struct.std::_Vector_base.120" = 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.156" }
%"class.std::shared_ptr.156" = type { %"class.std::__shared_ptr.157" }
%"class.std::__shared_ptr.157" = 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.sbo::small_vector" = type { %"class.std::vector.159" }
%"class.std::vector.159" = type { %"struct.std::_Vector_base.160" }
%"struct.std::_Vector_base.160" = 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.std::shared_ptr.153" = type { %"class.std::__shared_ptr.154" }
%"class.std::__shared_ptr.154" = type { ptr, %"class.std::__shared_count" }
%"class.hipsycl::rt::result" = type { %"class.std::unique_ptr" }
%"class.std::unique_ptr" = type { %"struct.std::__uniq_ptr_data" }
%"struct.std::__uniq_ptr_data" = type { %"class.std::__uniq_ptr_impl" }
%"class.std::__uniq_ptr_impl" = type { %"class.std::tuple" }
%"class.std::tuple" = type { %"struct.std::_Tuple_impl" }
%"struct.std::_Tuple_impl" = type { %"struct.std::_Head_base.67" }
%"struct.std::_Head_base.67" = type { ptr }

@.str = external hidden unnamed_addr constant [97 x i8], align 1
@.str.63 = external hidden unnamed_addr constant [2 x i8], align 1
@__PRETTY_FUNCTION__.pcudaStreamWaitEvent = external hidden unnamed_addr constant [97 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 noundef nonnull align 8 dereferenceable(104) ptr @_ZN7hipsycl2rt5pcuda17pcuda_application3getEv() local_unnamed_addr #1

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

declare noundef i32 @_ZN7hipsycl2rt5pcuda27get_most_recent_pcuda_errorEv() local_unnamed_addr #1

declare i32 @__gxx_personality_v0(...)

; Function Attrs: mustprogress uwtable
declare hidden noundef ptr @_ZN7hipsycl2rt5pcuda12_GLOBAL__N_122queue_or_default_queueEPNS1_6streamE(ptr noundef) local_unnamed_addr #2

declare noundef zeroext i1 @_ZNK7hipsycl2rt6result10is_successEv(ptr noundef nonnull align 8 dereferenceable(8)) local_unnamed_addr #1

declare void @_ZN7hipsycl2rt5pcuda20register_pcuda_errorERKNS0_6resultE10pcudaError(ptr noundef nonnull align 8 dereferenceable(8), i32 noundef) local_unnamed_addr #1

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

declare noundef nonnull align 8 dereferenceable(40) ptr @_ZN7hipsycl2rt5pcuda17pcuda_application8pcuda_rtEv(ptr noundef nonnull align 8 dereferenceable(104)) local_unnamed_addr #1

; Function Attrs: mustprogress uwtable
declare dso_local noundef ptr @_ZNK7hipsycl2rt5pcuda13pcuda_runtime6get_rtEv(ptr noundef nonnull align 8 dereferenceable(40)) local_unnamed_addr #2 align 2

; Function Attrs: mustprogress nounwind uwtable
declare dso_local noundef nonnull align 8 dereferenceable(72) ptr @_ZN7hipsycl2rt7runtime8backendsEv(ptr noundef nonnull align 8 dereferenceable(608)) local_unnamed_addr #4 align 2

declare noundef ptr @_ZNK7hipsycl2rt15backend_manager3getENS0_10backend_idE(ptr noundef nonnull align 8 dereferenceable(72), i32 noundef) local_unnamed_addr #1

declare noundef i32 @_ZNK7hipsycl2rt9device_id11get_backendEv(ptr noundef nonnull align 4 dereferenceable(16)) local_unnamed_addr #1

; Function Attrs: noreturn nounwind
declare void @__assert_fail(ptr noundef, ptr noundef, i32 noundef, ptr noundef) local_unnamed_addr #5

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

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

declare noundef zeroext i1 @_ZNK7hipsycl2rt5pcuda5event11is_recordedEv(ptr noundef nonnull align 8 dereferenceable(40)) local_unnamed_addr #1

; Function Attrs: mustprogress uwtable
define dso_local noundef i32 @pcudaStreamWaitEvent(ptr noundef %0, ptr noundef %1, i32 noundef %2) local_unnamed_addr #2 personality ptr @__gxx_personality_v0 {
  %4 = alloca %"class.std::shared_ptr.150", align 8
  %5 = alloca %"class.hipsycl::rt::execution_hints", align 8
  %6 = alloca %"class.sbo::small_vector", align 8
  %7 = alloca ptr, align 8
  %8 = alloca ptr, align 8
  %9 = alloca %"class.std::shared_ptr.153", align 8
  %10 = alloca %"class.hipsycl::rt::device_id", align 8
  %11 = alloca %"class.hipsycl::rt::device_id", align 8
  %12 = alloca %"class.hipsycl::rt::device_id", align 8
  %13 = alloca %"class.hipsycl::rt::device_id", align 8
  %14 = alloca %"class.hipsycl::rt::device_id", align 8
  %15 = alloca %"class.hipsycl::rt::result", align 8
  %16 = alloca %"class.hipsycl::rt::result", align 8
  %17 = alloca %"class.hipsycl::rt::result", align 8
  %18 = tail call noundef i32 @_ZN7hipsycl2rt5pcuda27get_most_recent_pcuda_errorEv()
  %.not = icmp eq i32 %18, 0
  br i1 %.not, label %19, label %177

19:                                               ; preds = %3
  %20 = tail call noundef ptr @_ZN7hipsycl2rt5pcuda12_GLOBAL__N_122queue_or_default_queueEPNS1_6streamE(ptr noundef %0)
  %.not42 = icmp eq ptr %20, null
  %.not43 = icmp eq ptr %1, null
  %or.cond = or i1 %.not43, %.not42
  br i1 %or.cond, label %177, label %21

21:                                               ; preds = %19
  %22 = tail call noundef zeroext i1 @_ZNK7hipsycl2rt5pcuda5event11is_recordedEv(ptr noundef nonnull align 8 dereferenceable(40) %1)
  br i1 %22, label %23, label %177

23:                                               ; preds = %21
  call void @llvm.lifetime.start.p0(i64 16, ptr nonnull %4) #8
  call void @llvm.lifetime.start.p0(i64 136, ptr nonnull %5) #8
  call void @llvm.memset.p0.i64(ptr noundef nonnull align 8 dereferenceable(136) %5, i8 0, i64 136, i1 false)
  call void @_ZN7hipsycl2rt15execution_hintsC2Ev(ptr noundef nonnull align 8 dereferenceable(132) %5) #8
  call void @llvm.lifetime.start.p0(i64 160, ptr nonnull %6) #8
  call void @_ZN3sbo12small_vectorISt10shared_ptrIN7hipsycl2rt8dag_nodeEELm8EEC2Ev(ptr noundef nonnull align 8 dereferenceable(160) %6) #8
  call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %7) #8
  store ptr null, ptr %7, align 8, !tbaa !5
  call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %8) #8
  %24 = invoke noundef nonnull align 8 dereferenceable(104) ptr @_ZN7hipsycl2rt5pcuda17pcuda_application3getEv()
          to label %25 unwind label %54

25:                                               ; preds = %23
  %26 = invoke noundef nonnull align 8 dereferenceable(40) ptr @_ZN7hipsycl2rt5pcuda17pcuda_application8pcuda_rtEv(ptr noundef nonnull align 8 dereferenceable(104) %24)
          to label %27 unwind label %54

27:                                               ; preds = %25
  %28 = invoke noundef ptr @_ZNK7hipsycl2rt5pcuda13pcuda_runtime6get_rtEv(ptr noundef nonnull align 8 dereferenceable(40) %26)
          to label %29 unwind label %54

29:                                               ; preds = %27
  store ptr %28, ptr %8, align 8, !tbaa !9
  invoke void @_ZSt11make_sharedIN7hipsycl2rt8dag_nodeEJNS1_15execution_hintsEN3sbo12small_vectorISt10shared_ptrIS2_ELm8EEEDnPNS1_7runtimeEEES6_INSt9enable_ifIXntsr8is_arrayIT_EE5valueESC_E4typeEEDpOT0_(ptr dead_on_unwind nonnull writable sret(%"class.std::shared_ptr.150") align 8 %4, ptr noundef nonnull align 8 dereferenceable(132) %5, ptr noundef nonnull align 8 dereferenceable(160) %6, ptr noundef nonnull align 8 dereferenceable(8) %7, ptr noundef nonnull align 8 dereferenceable(8) %8)
          to label %30 unwind label %54

30:                                               ; preds = %29
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %8) #8
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %7) #8
  call void @_ZNSt6vectorISt10shared_ptrIN7hipsycl2rt8dag_nodeEEN3sbo29small_buffer_vector_allocatorIS4_Lm8ES4_EEED2Ev(ptr noundef nonnull align 8 dereferenceable(160) %6) #8
  call void @llvm.lifetime.end.p0(i64 160, ptr nonnull %6) #8
  call void @_ZN7hipsycl2rt15execution_hintsD2Ev(ptr noundef nonnull align 8 dereferenceable(132) %5) #8
  call void @llvm.lifetime.end.p0(i64 136, ptr nonnull %5) #8
  %31 = call noundef ptr @_ZNKSt19__shared_ptr_accessIN7hipsycl2rt8dag_nodeELN9__gnu_cxx12_Lock_policyE2ELb0ELb0EEptEv(ptr noundef nonnull align 1 dereferenceable(1) %4) #8
  call void @_ZNK7hipsycl2rt5pcuda5event20get_event_shared_ptrEv(ptr dead_on_unwind nonnull writable sret(%"class.std::shared_ptr.153") align 8 %9, ptr noundef nonnull align 8 dereferenceable(40) %1)
  invoke void @_ZN7hipsycl2rt8dag_node14mark_submittedESt10shared_ptrINS0_14dag_node_eventEE(ptr noundef nonnull align 8 dereferenceable(384) %31, ptr noundef nonnull %9)
          to label %32 unwind label %58

32:                                               ; preds = %30
  call void @_ZNSt12__shared_ptrIN7hipsycl2rt14dag_node_eventELN9__gnu_cxx12_Lock_policyE2EED2Ev(ptr noundef nonnull align 8 dereferenceable(16) %9) #8
  %33 = call noundef ptr @_ZNKSt19__shared_ptr_accessIN7hipsycl2rt8dag_nodeELN9__gnu_cxx12_Lock_policyE2ELb0ELb0EEptEv(ptr noundef nonnull align 1 dereferenceable(1) %4) #8
  %34 = call { i64, i64 } @_ZNK7hipsycl2rt5pcuda5event10get_deviceEv(ptr noundef nonnull align 8 dereferenceable(40) %1)
  %35 = extractvalue { i64, i64 } %34, 0
  %36 = extractvalue { i64, i64 } %34, 1
  invoke void @_ZN7hipsycl2rt8dag_node16assign_to_deviceENS0_9device_idE(ptr noundef nonnull align 8 dereferenceable(384) %33, i64 %35, i64 %36)
          to label %37 unwind label %56

37:                                               ; preds = %32
  call void @llvm.lifetime.start.p0(i64 16, ptr nonnull %10) #8
  %38 = call { i64, i64 } @_ZNK7hipsycl2rt5pcuda5event10get_deviceEv(ptr noundef nonnull align 8 dereferenceable(40) %1)
  %39 = extractvalue { i64, i64 } %38, 0
  store i64 %39, ptr %10, align 8
  %40 = getelementptr inbounds i8, ptr %10, i64 8
  %41 = extractvalue { i64, i64 } %38, 1
  store i64 %41, ptr %40, align 8
  %42 = invoke noundef i32 @_ZNK7hipsycl2rt9device_id11get_backendEv(ptr noundef nonnull align 4 dereferenceable(16) %10)
          to label %43 unwind label %60

43:                                               ; preds = %37
  call void @llvm.lifetime.start.p0(i64 16, ptr nonnull %11) #8
  %44 = load ptr, ptr %20, align 8, !tbaa !11
  %45 = getelementptr inbounds i8, ptr %44, i64 72
  %46 = load ptr, ptr %45, align 8
  %47 = invoke { i64, i64 } %46(ptr noundef nonnull align 8 dereferenceable(8) %20)
          to label %48 unwind label %62

48:                                               ; preds = %43
  %49 = extractvalue { i64, i64 } %47, 0
  store i64 %49, ptr %11, align 8
  %50 = getelementptr inbounds i8, ptr %11, i64 8
  %51 = extractvalue { i64, i64 } %47, 1
  store i64 %51, ptr %50, align 8
  %52 = invoke noundef i32 @_ZNK7hipsycl2rt9device_id11get_backendEv(ptr noundef nonnull align 4 dereferenceable(16) %11)
          to label %53 unwind label %62

53:                                               ; preds = %48
  %.not45 = icmp eq i32 %42, %52
  call void @llvm.lifetime.end.p0(i64 16, ptr nonnull %11) #8
  call void @llvm.lifetime.end.p0(i64 16, ptr nonnull %10) #8
  br i1 %.not45, label %65, label %.critedge

54:                                               ; preds = %29, %27, %25, %23
  %55 = landingpad { ptr, i32 }
          cleanup
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %8) #8
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %7) #8
  call void @_ZNSt6vectorISt10shared_ptrIN7hipsycl2rt8dag_nodeEEN3sbo29small_buffer_vector_allocatorIS4_Lm8ES4_EEED2Ev(ptr noundef nonnull align 8 dereferenceable(160) %6) #8
  call void @llvm.lifetime.end.p0(i64 160, ptr nonnull %6) #8
  call void @_ZN7hipsycl2rt15execution_hintsD2Ev(ptr noundef nonnull align 8 dereferenceable(132) %5) #8
  call void @llvm.lifetime.end.p0(i64 136, ptr nonnull %5) #8
  br label %176

56:                                               ; preds = %32
  %57 = landingpad { ptr, i32 }
          cleanup
  br label %175

58:                                               ; preds = %30
  %59 = landingpad { ptr, i32 }
          cleanup
  call void @_ZNSt12__shared_ptrIN7hipsycl2rt14dag_node_eventELN9__gnu_cxx12_Lock_policyE2EED2Ev(ptr noundef nonnull align 8 dereferenceable(16) %9) #8
  br label %175

60:                                               ; preds = %37
  %61 = landingpad { ptr, i32 }
          cleanup
  br label %64

62:                                               ; preds = %48, %43
  %63 = landingpad { ptr, i32 }
          cleanup
  call void @llvm.lifetime.end.p0(i64 16, ptr nonnull %11) #8
  br label %64

64:                                               ; preds = %62, %60
  %.pn = phi { ptr, i32 } [ %63, %62 ], [ %61, %60 ]
  call void @llvm.lifetime.end.p0(i64 16, ptr nonnull %10) #8
  br label %175

65:                                               ; preds = %53
  %66 = invoke noundef nonnull align 8 dereferenceable(104) ptr @_ZN7hipsycl2rt5pcuda17pcuda_application3getEv()
          to label %67 unwind label %82

67:                                               ; preds = %65
  %68 = invoke noundef nonnull align 8 dereferenceable(40) ptr @_ZN7hipsycl2rt5pcuda17pcuda_application8pcuda_rtEv(ptr noundef nonnull align 8 dereferenceable(104) %66)
          to label %69 unwind label %82

69:                                               ; preds = %67
  %70 = invoke noundef ptr @_ZNK7hipsycl2rt5pcuda13pcuda_runtime6get_rtEv(ptr noundef nonnull align 8 dereferenceable(40) %68)
          to label %71 unwind label %82

71:                                               ; preds = %69
  %72 = call noundef nonnull align 8 dereferenceable(72) ptr @_ZN7hipsycl2rt7runtime8backendsEv(ptr noundef nonnull align 8 dereferenceable(608) %70)
  call void @llvm.lifetime.start.p0(i64 16, ptr nonnull %12) #8
  %73 = call { i64, i64 } @_ZNK7hipsycl2rt5pcuda5event10get_deviceEv(ptr noundef nonnull align 8 dereferenceable(40) %1)
  %74 = extractvalue { i64, i64 } %73, 0
  store i64 %74, ptr %12, align 8
  %75 = getelementptr inbounds i8, ptr %12, i64 8
  %76 = extractvalue { i64, i64 } %73, 1
  store i64 %76, ptr %75, align 8
  %77 = invoke noundef i32 @_ZNK7hipsycl2rt9device_id11get_backendEv(ptr noundef nonnull align 4 dereferenceable(16) %12)
          to label %78 unwind label %84

78:                                               ; preds = %71
  %79 = invoke noundef ptr @_ZNK7hipsycl2rt15backend_manager3getENS0_10backend_idE(ptr noundef nonnull align 8 dereferenceable(72) %72, i32 noundef %77)
          to label %80 unwind label %84

80:                                               ; preds = %78
  call void @llvm.lifetime.end.p0(i64 16, ptr nonnull %12) #8
  %.not46 = icmp eq ptr %79, null
  br i1 %.not46, label %81, label %86

81:                                               ; preds = %80
  call void @__assert_fail(ptr noundef nonnull @.str.63, ptr noundef nonnull @.str, i32 noundef 849, ptr noundef nonnull @__PRETTY_FUNCTION__.pcudaStreamWaitEvent) #9
  unreachable

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

84:                                               ; preds = %78, %71
  %85 = landingpad { ptr, i32 }
          cleanup
  call void @llvm.lifetime.end.p0(i64 16, ptr nonnull %12) #8
  br label %175

86:                                               ; preds = %80
  %87 = load ptr, ptr %79, align 8, !tbaa !11
  %88 = getelementptr inbounds i8, ptr %87, i64 24
  %89 = load ptr, ptr %88, align 8
  %90 = invoke noundef ptr %89(ptr noundef nonnull align 8 dereferenceable(8) %79)
          to label %91 unwind label %135

91:                                               ; preds = %86
  call void @llvm.lifetime.start.p0(i64 16, ptr nonnull %13) #8
  %92 = call { i64, i64 } @_ZNK7hipsycl2rt5pcuda5event10get_deviceEv(ptr noundef nonnull align 8 dereferenceable(40) %1)
  %93 = extractvalue { i64, i64 } %92, 0
  store i64 %93, ptr %13, align 8
  %94 = getelementptr inbounds i8, ptr %13, i64 8
  %95 = extractvalue { i64, i64 } %92, 1
  store i64 %95, ptr %94, align 8
  %96 = invoke noundef i32 @_ZNK7hipsycl2rt9device_id6get_idEv(ptr noundef nonnull align 4 dereferenceable(16) %13)
          to label %97 unwind label %137

97:                                               ; preds = %91
  %98 = sext i32 %96 to i64
  %99 = load ptr, ptr %90, align 8, !tbaa !11
  %100 = getelementptr inbounds i8, ptr %99, i64 16
  %101 = load ptr, ptr %100, align 8
  %102 = invoke noundef ptr %101(ptr noundef nonnull align 8 dereferenceable(8) %90, i64 noundef %98)
          to label %103 unwind label %137

103:                                              ; preds = %97
  %104 = load ptr, ptr %102, align 8, !tbaa !11
  %105 = getelementptr inbounds i8, ptr %104, i64 96
  %106 = load ptr, ptr %105, align 8
  %107 = invoke noundef i64 %106(ptr noundef nonnull align 8 dereferenceable(8) %102)
          to label %108 unwind label %137

108:                                              ; preds = %103
  call void @llvm.lifetime.end.p0(i64 16, ptr nonnull %13) #8
  %109 = load ptr, ptr %79, align 8, !tbaa !11
  %110 = getelementptr inbounds i8, ptr %109, i64 24
  %111 = load ptr, ptr %110, align 8
  %112 = invoke noundef ptr %111(ptr noundef nonnull align 8 dereferenceable(8) %79)
          to label %113 unwind label %139

113:                                              ; preds = %108
  call void @llvm.lifetime.start.p0(i64 16, ptr nonnull %14) #8
  %114 = load ptr, ptr %20, align 8, !tbaa !11
  %115 = getelementptr inbounds i8, ptr %114, i64 72
  %116 = load ptr, ptr %115, align 8
  %117 = invoke { i64, i64 } %116(ptr noundef nonnull align 8 dereferenceable(8) %20)
          to label %118 unwind label %141

118:                                              ; preds = %113
  %119 = extractvalue { i64, i64 } %117, 0
  store i64 %119, ptr %14, align 8
  %120 = getelementptr inbounds i8, ptr %14, i64 8
  %121 = extractvalue { i64, i64 } %117, 1
  store i64 %121, ptr %120, align 8
  %122 = invoke noundef i32 @_ZNK7hipsycl2rt9device_id6get_idEv(ptr noundef nonnull align 4 dereferenceable(16) %14)
          to label %123 unwind label %141

123:                                              ; preds = %118
  %124 = sext i32 %122 to i64
  %125 = load ptr, ptr %112, align 8, !tbaa !11
  %126 = getelementptr inbounds i8, ptr %125, i64 16
  %127 = load ptr, ptr %126, align 8
  %128 = invoke noundef ptr %127(ptr noundef nonnull align 8 dereferenceable(8) %112, i64 noundef %124)
          to label %129 unwind label %141

129:                                              ; preds = %123
  %130 = load ptr, ptr %128, align 8, !tbaa !11
  %131 = getelementptr inbounds i8, ptr %130, i64 96
  %132 = load ptr, ptr %131, align 8
  %133 = invoke noundef i64 %132(ptr noundef nonnull align 8 dereferenceable(8) %128)
          to label %134 unwind label %141

134:                                              ; preds = %129
  call void @llvm.lifetime.end.p0(i64 16, ptr nonnull %14) #8
  %.not51 = icmp eq i64 %107, %133
  call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %15) #8
  call void @_ZN7hipsycl2rt6resultC2Ev(ptr noundef nonnull align 8 dereferenceable(8) %15) #8
  br i1 %.not51, label %143, label %155

135:                                              ; preds = %86
  %136 = landingpad { ptr, i32 }
          cleanup
  br label %175

137:                                              ; preds = %103, %97, %91
  %138 = landingpad { ptr, i32 }
          cleanup
  call void @llvm.lifetime.end.p0(i64 16, ptr nonnull %13) #8
  br label %175

139:                                              ; preds = %108
  %140 = landingpad { ptr, i32 }
          cleanup
  br label %175

141:                                              ; preds = %129, %123, %118, %113
  %142 = landingpad { ptr, i32 }
          cleanup
  call void @llvm.lifetime.end.p0(i64 16, ptr nonnull %14) #8
  br label %175

143:                                              ; preds = %134
  call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %16) #8
  %144 = load ptr, ptr %20, align 8, !tbaa !11
  %145 = getelementptr inbounds i8, ptr %144, i64 48
  %146 = load ptr, ptr %145, align 8
  invoke void %146(ptr dead_on_unwind nonnull writable sret(%"class.hipsycl::rt::result") align 8 %16, ptr noundef nonnull align 8 dereferenceable(8) %20, ptr noundef nonnull align 8 dereferenceable(16) %4)
          to label %147 unwind label %150

147:                                              ; preds = %143
  %148 = invoke noundef nonnull align 8 dereferenceable(8) ptr @_ZN7hipsycl2rt6resultaSEOS1_(ptr noundef nonnull align 8 dereferenceable(8) %15, ptr noundef nonnull align 8 dereferenceable(8) %16)
          to label %149 unwind label %152

149:                                              ; preds = %147
  call void @_ZN7hipsycl2rt6resultD2Ev(ptr noundef nonnull align 8 dereferenceable(8) %16) #8
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %16) #8
  br label %167

150:                                              ; preds = %143
  %151 = landingpad { ptr, i32 }
          cleanup
  br label %154

152:                                              ; preds = %147
  %153 = landingpad { ptr, i32 }
          cleanup
  call void @_ZN7hipsycl2rt6resultD2Ev(ptr noundef nonnull align 8 dereferenceable(8) %16) #8
  br label %154

154:                                              ; preds = %152, %150
  %.pn54 = phi { ptr, i32 } [ %153, %152 ], [ %151, %150 ]
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %16) #8
  br label %174

.critedge:                                        ; preds = %53
  call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %15) #8
  call void @_ZN7hipsycl2rt6resultC2Ev(ptr noundef nonnull align 8 dereferenceable(8) %15) #8
  br label %155

155:                                              ; preds = %.critedge, %134
  call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %17) #8
  %156 = load ptr, ptr %20, align 8, !tbaa !11
  %157 = getelementptr inbounds i8, ptr %156, i64 56
  %158 = load ptr, ptr %157, align 8
  invoke void %158(ptr dead_on_unwind nonnull writable sret(%"class.hipsycl::rt::result") align 8 %17, ptr noundef nonnull align 8 dereferenceable(8) %20, ptr noundef nonnull align 8 dereferenceable(16) %4)
          to label %159 unwind label %162

159:                                              ; preds = %155
  %160 = invoke noundef nonnull align 8 dereferenceable(8) ptr @_ZN7hipsycl2rt6resultaSEOS1_(ptr noundef nonnull align 8 dereferenceable(8) %15, ptr noundef nonnull align 8 dereferenceable(8) %17)
          to label %161 unwind label %164

161:                                              ; preds = %159
  call void @_ZN7hipsycl2rt6resultD2Ev(ptr noundef nonnull align 8 dereferenceable(8) %17) #8
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %17) #8
  br label %167

162:                                              ; preds = %155
  %163 = landingpad { ptr, i32 }
          cleanup
  br label %166

164:                                              ; preds = %159
  %165 = landingpad { ptr, i32 }
          cleanup
  call void @_ZN7hipsycl2rt6resultD2Ev(ptr noundef nonnull align 8 dereferenceable(8) %17) #8
  br label %166

166:                                              ; preds = %164, %162
  %.pn52 = phi { ptr, i32 } [ %165, %164 ], [ %163, %162 ]
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %17) #8
  br label %174

167:                                              ; preds = %161, %149
  %168 = invoke noundef zeroext i1 @_ZNK7hipsycl2rt6result10is_successEv(ptr noundef nonnull align 8 dereferenceable(8) %15)
          to label %169 unwind label %171

169:                                              ; preds = %167
  br i1 %168, label %173, label %170

170:                                              ; preds = %169
  invoke void @_ZN7hipsycl2rt5pcuda20register_pcuda_errorERKNS0_6resultE10pcudaError(ptr noundef nonnull align 8 dereferenceable(8) %15, i32 noundef 30)
          to label %173 unwind label %171

171:                                              ; preds = %170, %167
  %172 = landingpad { ptr, i32 }
          cleanup
  br label %174

173:                                              ; preds = %169, %170
  %.0 = phi i32 [ 30, %170 ], [ 0, %169 ]
  call void @_ZN7hipsycl2rt6resultD2Ev(ptr noundef nonnull align 8 dereferenceable(8) %15) #8
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %15) #8
  call void @_ZNSt12__shared_ptrIN7hipsycl2rt8dag_nodeELN9__gnu_cxx12_Lock_policyE2EED2Ev(ptr noundef nonnull align 8 dereferenceable(16) %4) #8
  call void @llvm.lifetime.end.p0(i64 16, ptr nonnull %4) #8
  br label %177

174:                                              ; preds = %171, %166, %154
  %.pn56 = phi { ptr, i32 } [ %172, %171 ], [ %.pn54, %154 ], [ %.pn52, %166 ]
  call void @_ZN7hipsycl2rt6resultD2Ev(ptr noundef nonnull align 8 dereferenceable(8) %15) #8
  call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %15) #8
  br label %175

175:                                              ; preds = %64, %174, %135, %137, %141, %139, %84, %82, %58, %56
  %.pn56.pn.pn = phi { ptr, i32 } [ %57, %56 ], [ %59, %58 ], [ %.pn56, %174 ], [ %.pn, %64 ], [ %85, %84 ], [ %83, %82 ], [ %138, %137 ], [ %136, %135 ], [ %142, %141 ], [ %140, %139 ]
  call void @_ZNSt12__shared_ptrIN7hipsycl2rt8dag_nodeELN9__gnu_cxx12_Lock_policyE2EED2Ev(ptr noundef nonnull align 8 dereferenceable(16) %4) #8
  br label %176

176:                                              ; preds = %175, %54
  %.pn56.pn.pn.pn = phi { ptr, i32 } [ %.pn56.pn.pn, %175 ], [ %55, %54 ]
  call void @llvm.lifetime.end.p0(i64 16, ptr nonnull %4) #8
  resume { ptr, i32 } %.pn56.pn.pn.pn

177:                                              ; preds = %173, %19, %21, %3
  %.2 = phi i32 [ %18, %3 ], [ %.0, %173 ], [ 33, %19 ], [ 0, %21 ]
  ret i32 %.2
}

; Function Attrs: inlinehint mustprogress uwtable
declare dso_local void @_ZSt11make_sharedIN7hipsycl2rt8dag_nodeEJNS1_15execution_hintsEN3sbo12small_vectorISt10shared_ptrIS2_ELm8EEEDnPNS1_7runtimeEEES6_INSt9enable_ifIXntsr8is_arrayIT_EE5valueESC_E4typeEEDpOT0_(ptr dead_on_unwind noalias writable sret(%"class.std::shared_ptr.150") align 8, ptr noundef nonnull align 8 dereferenceable(132), ptr noundef nonnull align 8 dereferenceable(160), ptr noundef nonnull align 8 dereferenceable(8), ptr noundef nonnull align 8 dereferenceable(8)) local_unnamed_addr #7

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

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

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

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

declare void @_ZN7hipsycl2rt8dag_node14mark_submittedESt10shared_ptrINS0_14dag_node_eventEE(ptr noundef nonnull align 8 dereferenceable(384), ptr noundef) local_unnamed_addr #1

; Function Attrs: mustprogress nounwind uwtable
declare dso_local void @_ZNK7hipsycl2rt5pcuda5event20get_event_shared_ptrEv(ptr dead_on_unwind noalias writable sret(%"class.std::shared_ptr.153") align 8, ptr noundef nonnull align 8 dereferenceable(40)) local_unnamed_addr #4 align 2

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

declare void @_ZN7hipsycl2rt8dag_node16assign_to_deviceENS0_9device_idE(ptr noundef nonnull align 8 dereferenceable(384), i64, i64) local_unnamed_addr #1

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

declare noundef i32 @_ZNK7hipsycl2rt9device_id6get_idEv(ptr noundef nonnull align 4 dereferenceable(16)) local_unnamed_addr #1

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

declare noundef nonnull align 8 dereferenceable(8) ptr @_ZN7hipsycl2rt6resultaSEOS1_(ptr noundef nonnull align 8 dereferenceable(8), ptr noundef nonnull align 8 dereferenceable(8)) local_unnamed_addr #1

attributes #0 = { mustprogress nocallback nofree nosync nounwind willreturn memory(argmem: readwrite) }
attributes #1 = { "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cmov,+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #2 = { mustprogress uwtable "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cmov,+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #3 = { inlinehint mustprogress nounwind uwtable "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cmov,+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #4 = { mustprogress nounwind uwtable "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cmov,+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "tune-cpu"="generic" }
attributes #5 = { noreturn nounwind "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 = { mustprogress nocallback nofree nounwind willreturn memory(argmem: write) }
attributes #7 = { 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 #8 = { nounwind }
attributes #9 = { noreturn 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 = !{!"std::nullptr_t", !7, i64 0}
!7 = !{!"omnipotent char", !8, i64 0}
!8 = !{!"Simple C++ TBAA"}
!9 = !{!10, !10, i64 0}
!10 = !{!"any pointer", !7, i64 0}
!11 = !{!12, !12, i64 0}
!12 = !{!"vtable pointer", !8, i64 0}


; Effective O3 passes: EarlyCSEPass,GlobalOptPass,IPSCCPPass,InferAlignmentPass,InferFunctionAttrsPass,InstCombinePass,PostOrderFunctionAttrsPass,ReassociatePass,SROAPass,SimplifyCFGPass,TailCallElimPass