From 811db848c397c4e2403e37f2b02fc28e167b56d0 Mon Sep 17 00:00:00 2001
From: Artur Gainullin <artur.gainullin@intel.com>
Date: Tue, 10 Sep 2024 18:22:14 -0700
Subject: [PATCH] [SYCL] Fix handling of interop events for barrier with
 waitlist (#15352)

Currently Command::getUrEventsBlocking is responsible for preparing a
waitlist of UR events for the barrier.
This method used wrong assumption that if isEnqueued() returns false for
the event then it doesn't have UR handle because it was not enqueued. So
if there is an associated command we would enqueue it to get the desired
UR handle, or we would just ignore this event if there is no associated
command.

Problem is that sycl::event created with interoperability constructor
has isEnqueued() as false (as it is not enqueued by SYCL RT) but it has
UR handle provided by user.
Before this patch we just ignored such event as it doesn't have
associated command and we didn't put it to the resulting list.

This patch fixes this problem by handling interop events properly in
this code path.
---
 sycl/source/detail/event_impl.hpp             |  7 +++
 sycl/source/detail/scheduler/commands.cpp     |  9 ++--
 .../barrier_waitlist_with_interop_event.cpp   | 47 +++++++++++++++++++
 3 files changed, 60 insertions(+), 3 deletions(-)
 create mode 100644 sycl/test-e2e/Regression/barrier_waitlist_with_interop_event.cpp

diff --git a/sycl/source/detail/event_impl.hpp b/sycl/source/detail/event_impl.hpp
index b560d721728a6..312bb589760b7 100644
--- a/sycl/source/detail/event_impl.hpp
+++ b/sycl/source/detail/event_impl.hpp
@@ -329,6 +329,13 @@ class event_impl {
 
   bool isProfilingTagEvent() const noexcept { return MProfilingTagEvent; }
 
+  // Check if this event is an interoperability event.
+  bool isInterop() const noexcept {
+    // As an indication of interoperability event, we use the absence of the
+    // queue and command, as well as the fact that it is not in enqueued state.
+    return MEvent && MQueue.expired() && !MIsEnqueued && !MCommand;
+  }
+
 protected:
   // When instrumentation is enabled emits trace event for event wait begin and
   // returns the telemetry event generated for the wait
diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp
index 953ad2bee0444..c5e8fc2c3a2cd 100644
--- a/sycl/source/detail/scheduler/commands.cpp
+++ b/sycl/source/detail/scheduler/commands.cpp
@@ -265,9 +265,12 @@ std::vector<ur_event_handle_t> Command::getUrEventsBlocking(
     if (EventImpl->isDefaultConstructed() || EventImpl->isHost() ||
         EventImpl->isNOP())
       continue;
-    // In this path nullptr native event means that the command has not been
-    // enqueued. It may happen if async enqueue in a host task is involved.
-    if (!EventImpl->isEnqueued()) {
+
+    // If command has not been enqueued then we have to enqueue it.
+    // It may happen if async enqueue in a host task is involved.
+    // Interoperability events are special cases and they are not enqueued, as
+    // they don't have an associated queue and command.
+    if (!EventImpl->isInterop() && !EventImpl->isEnqueued()) {
       if (!EventImpl->getCommand() ||
           !static_cast<Command *>(EventImpl->getCommand())->producesPiEvent())
         continue;
diff --git a/sycl/test-e2e/Regression/barrier_waitlist_with_interop_event.cpp b/sycl/test-e2e/Regression/barrier_waitlist_with_interop_event.cpp
new file mode 100644
index 0000000000000..f5a54f1a67dc2
--- /dev/null
+++ b/sycl/test-e2e/Regression/barrier_waitlist_with_interop_event.cpp
@@ -0,0 +1,47 @@
+// REQUIRES: level_zero, level_zero_dev_kit
+// RUN: %{build} %level_zero_options -o %t.out
+// RUN: %{run} %t.out
+// UNSUPPORTED: ze_debug
+
+#include <level_zero/ze_api.h>
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/backend/level_zero.hpp>
+#include <sycl/usm.hpp>
+
+// Test checks the case when an interoperability event is passed as a dependency
+// to the barrier. In such case, waiting for the event produced by barrier must
+// guarantee completion of the interoperability event.
+
+using namespace sycl;
+
+int main() {
+  sycl::queue Queue;
+  if (!Queue.get_device().get_info<info::device::usm_shared_allocations>())
+    return 0;
+
+  const size_t N = 1024;
+  int *Data = sycl::malloc_shared<int>(N, Queue);
+  auto FillEvent = Queue.fill(Data, 0, N);
+  auto FillZeEvent = get_native<backend::ext_oneapi_level_zero>(FillEvent);
+
+  backend_input_t<backend::ext_oneapi_level_zero, event> EventInteropInput = {
+      FillZeEvent};
+  EventInteropInput.Ownership = sycl::ext::oneapi::level_zero::ownership::keep;
+  auto EventInterop = make_event<backend::ext_oneapi_level_zero>(
+      EventInteropInput, Queue.get_context());
+
+  auto BarrierEvent = Queue.ext_oneapi_submit_barrier({EventInterop});
+  BarrierEvent.wait();
+
+  if (EventInterop.get_info<sycl::info::event::command_execution_status>() !=
+      sycl::info::event_command_status::complete) {
+    Queue.wait();
+    sycl::free(Data, Queue);
+    return -1;
+  }
+
+  // Free the USM memory
+  sycl::free(Data, Queue);
+
+  return 0;
+}