Ht user events 1

sycl/source/detail/event_impl.cpp

@@ -54,7 +54,7 @@ event_impl::~event_impl() {
 
 void event_impl::waitInternal(bool *Success) {
   auto Handle = this->getHandle();
-  if (!MIsHostEvent && Handle) {
+  if (Handle) {
     // Wait for the native event
     ur_result_t Err =
         getAdapter()->call_nocheck<UrApiKind::urEventWait>(1, &Handle);
@@ -92,7 +92,7 @@ void event_impl::waitInternal(bool *Success) {
 }
 
 void event_impl::setComplete() {
-  if (MIsHostEvent || !this->getHandle()) {
+  if (!this->getHandle()) {
     {
       std::unique_lock<std::mutex> lock(MMutex);
 #ifndef NDEBUG
@@ -138,7 +138,6 @@ const AdapterPtr &event_impl::getAdapter() {
 void event_impl::setStateIncomplete() { MState = HES_NotComplete; }
 
 void event_impl::setContextImpl(const ContextImplPtr &Context) {
-  MIsHostEvent = Context == nullptr;
   MContext = Context;
 }
 
@@ -159,6 +158,14 @@ event_impl::event_impl(ur_event_handle_t Event, const context &SyclContext)
   }
 }
 
+void event_impl::allocateHostProfilingInfo() {
+  MHostProfilingInfo.reset(new HostProfilingInfo());
+  if (!MHostProfilingInfo)
+    throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
+                          "Out of host memory " +
+                              codeToString(UR_RESULT_ERROR_OUT_OF_HOST_MEMORY));
+}
+
 event_impl::event_impl(const QueueImplPtr &Queue)
     : MQueue{Queue}, MIsProfilingEnabled{!Queue || Queue->MIsProfilingEnabled},
       MFallbackProfiling{MIsProfilingEnabled && Queue &&
@@ -167,12 +174,7 @@ event_impl::event_impl(const QueueImplPtr &Queue)
     this->setContextImpl(Queue->getContextImplPtr());
   else {
     MState.store(HES_NotComplete);
-    MHostProfilingInfo.reset(new HostProfilingInfo());
-    if (!MHostProfilingInfo)
-      throw sycl::exception(
-          sycl::make_error_code(sycl::errc::runtime),
-          "Out of host memory " +
-              codeToString(UR_RESULT_ERROR_OUT_OF_HOST_MEMORY));
+    allocateHostProfilingInfo();
     return;
   }
   MState.store(HES_Complete);
@@ -400,7 +402,7 @@ uint64_t event_impl::get_profiling_info<info::event_profiling::command_end>() {
 
 template <> uint32_t event_impl::get_info<info::event::reference_count>() {
   auto Handle = this->getHandle();
-  if (!MIsHostEvent && Handle) {
+  if (Handle) {
     return get_event_info<info::event::reference_count>(Handle,
                                                         this->getAdapter());
   }
@@ -496,18 +498,20 @@ void HostProfilingInfo::start() { StartTime = getTimestamp(); }
 void HostProfilingInfo::end() { EndTime = getTimestamp(); }
 
 ur_native_handle_t event_impl::getNative() {
-  if (isHost())
+  auto Handle = getHandle();
+  if (MIsHostEvent && !Handle)
     return {};
-  initContextIfNeeded();
 
+  initContextIfNeeded();
   auto Adapter = getAdapter();
-  auto Handle = getHandle();
+
   if (MIsDefaultConstructed && !Handle) {
     auto TempContext = MContext.get()->getHandleRef();
     ur_event_native_properties_t NativeProperties{};
     ur_event_handle_t UREvent = nullptr;
     Adapter->call<UrApiKind::urEventCreateWithNativeHandle>(
         0, TempContext, &NativeProperties, &UREvent);
+    Adapter->call<UrApiKind::urEventHostSignal>(UREvent);
     this->setHandle(UREvent);
     Handle = UREvent;
   }
@@ -627,8 +631,14 @@ bool event_impl::isCompleted() {
 void event_impl::setCommand(void *Cmd) {
   MCommand = Cmd;
   auto TypedCommand = static_cast<Command *>(Cmd);
-  if (TypedCommand)
-    MIsHostEvent = TypedCommand->getWorkerContext() == nullptr;
+  if (TypedCommand && TypedCommand->getWorkerContext() == nullptr)
+    markAsHost();
+}
+
+void event_impl::markAsHost() {
+  MIsHostEvent = true;
+  if (!MHostProfilingInfo)
+    allocateHostProfilingInfo();
 }
 
 } // namespace detail

sycl/source/detail/event_impl.hpp

@@ -349,6 +349,8 @@ class event_impl {
     return MEvent && MQueue.expired() && !MIsEnqueued && !MCommand;
   }
 
+  void markAsHost();
+
 protected:
   // When instrumentation is enabled emits trace event for event wait begin and
   // returns the telemetry event generated for the wait
@@ -358,6 +360,7 @@ class event_impl {
   void instrumentationEpilog(void *TelementryEvent, const std::string &Name,
                              int32_t StreamID, uint64_t IId) const;
   void checkProfilingPreconditions() const;
+  void allocateHostProfilingInfo();
 
   std::atomic<ur_event_handle_t> MEvent = nullptr;
   // Stores submission time of command associated with event

sycl/source/detail/queue_impl.cpp

@@ -315,8 +315,10 @@ void queue_impl::addEvent(const event &Event) {
   }
   // As long as the queue supports urQueueFinish we only need to store events
   // for undiscarded, unenqueued commands and host tasks.
-  else if (MEmulateOOO ||
-           (EImpl->getHandle() == nullptr && !EImpl->isDiscarded())) {
+  // Event->isHost can be false for L0 on user events but we still need to sync
+  // host tasks explicitly.
+  else if (MEmulateOOO || ((EImpl->getHandle() == nullptr || EImpl->isHost()) &&
+                           !EImpl->isDiscarded())) {
     std::weak_ptr<event_impl> EventWeakPtr{EImpl};
     std::lock_guard<std::mutex> Lock{MMutex};
     MEventsWeak.push_back(std::move(EventWeakPtr));
@@ -637,7 +639,8 @@ void queue_impl::wait(const detail::code_location &CodeLoc) {
             EventImplWeakPtrIt->lock()) {
       // A nullptr UR event indicates that urQueueFinish will not cover it,
       // either because it's a host task event or an unenqueued one.
-      if (!SupportsPiFinish || nullptr == EventImplSharedPtr->getHandle()) {
+      if (!SupportsPiFinish || nullptr == EventImplSharedPtr->getHandle() ||
+          EventImplSharedPtr->isHost()) {
         EventImplSharedPtr->wait(EventImplSharedPtr);
       }
     }

sycl/source/detail/queue_impl.hpp

@@ -755,6 +755,11 @@ class queue_impl {
     return ResEvent;
   }
 
+  bool nativeHostTaskHandling() {
+    return std::getenv("SYCL_ENABLE_USER_EVENTS_PATH") &&
+           (MDevice->getBackend() == backend::ext_oneapi_level_zero);
+  }
+
 #ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   // CMPLRLLVM-66082
   // These methods are for accessing a member that should live in the