[SYCL] Reusable events extension implementation

llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td

@@ -106,6 +106,7 @@ def Aspectext_intel_eus_per_xe_core : Aspect<"ext_intel_eus_per_xe_core">;
 def Aspectext_intel_max_lanes_per_hw_thread : Aspect<"ext_intel_max_lanes_per_hw_thread">;
 def Aspectext_oneapi_ipc_physical_memory : Aspect<"ext_oneapi_ipc_physical_memory">;
 def Aspectext_oneapi_register_host_memory : Aspect<"ext_oneapi_register_host_memory">;
+def Aspectext_oneapi_per_event_profiling : Aspect<"ext_oneapi_per_event_profiling">;
 
 // Deprecated aspects
 def AspectInt64_base_atomics : Aspect<"int64_base_atomics">;
@@ -193,7 +194,8 @@ def : TargetInfo<"__TestAspectList",
     Aspectext_intel_eus_per_xe_core,
     Aspectext_intel_max_lanes_per_hw_thread,
     Aspectext_oneapi_ipc_physical_memory,
-    Aspectext_oneapi_register_host_memory],
+    Aspectext_oneapi_register_host_memory,
+    Aspectext_oneapi_per_event_profiling],
     []>;
 // This definition serves the only purpose of testing whether the deprecated aspect list defined in here and in SYCL RT
 // match.

sycl/include/sycl/ext/oneapi/experimental/reusable_events.hpp

@@ -0,0 +1,67 @@
+//==------- reusable_events.hpp --- SYCL reusable events -------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <sycl/context.hpp>
+#include <sycl/device.hpp>
+#include <sycl/event.hpp>
+#include <sycl/ext/oneapi/properties/properties.hpp>
+#include <sycl/platform.hpp>
+#include <sycl/queue.hpp>
+
+namespace sycl {
+inline namespace _V1 {
+namespace ext::oneapi::experimental {
+
+struct enable_profiling_key
+    : detail::compile_time_property_key<detail::PropKind::EnableProfiling> {
+  using value_t = property_value<enable_profiling_key>;
+};
+
+inline constexpr enable_profiling_key::value_t enable_profiling;
+
+namespace detail {
+
+enum make_event_flags : uint32_t { make_event_flag_enable_profiling = 1u };
+
+__SYCL_EXPORT sycl::event make_event(const sycl::context &ctxt, uint32_t Flags);
+
+template <typename PropertyListT> uint32_t getMakeEventFlags() {
+  uint32_t Flags = 0;
+  if constexpr (PropertyListT::template has_property<enable_profiling_key>())
+    Flags |= make_event_flag_enable_profiling;
+  return Flags;
+}
+} // namespace detail
+
+template <typename PropertyListT = empty_properties_t>
+inline sycl::event make_event(const sycl::context &ctxt,
+                              PropertyListT props = {}) {
+  static_assert(is_property_list_v<PropertyListT>,
+                "Props must be a sycl::ext::oneapi::experimental::properties");
+  (void)props;
+
+  return detail::make_event(ctxt, detail::getMakeEventFlags<PropertyListT>());
+}
+
+template <typename PropertyListT = empty_properties_t>
+inline sycl::event make_event(PropertyListT props = {}) {
+  sycl::device Dev;
+  sycl::context Ctx = Dev.get_platform().khr_get_default_context();
+  return make_event(Ctx, props);
+}
+
+__SYCL_EXPORT void enqueue_wait_event(sycl::queue q, const event &evt);
+__SYCL_EXPORT void enqueue_wait_events(sycl::queue q,
+                                       const std::vector<event> &evts);
+__SYCL_EXPORT void enqueue_signal_event(sycl::queue q, event &evt);
+
+} // namespace ext::oneapi::experimental
+} // namespace _V1
+} // namespace sycl

sycl/include/sycl/ext/oneapi/properties/property.hpp

@@ -194,8 +194,9 @@ enum PropKind : uint32_t {
   FastLink = 49,
   PhysicalMemoryEnableIPC = 50,
   RegisterHostMemoryReadOnly = 51,
+  EnableProfiling = 52,
   // PropKindSize must always be the last value.
-  PropKindSize = 52,
+  PropKindSize = 53,
 };
 
 template <typename PropertyT> struct PropertyToKind {

sycl/include/sycl/info/aspects.def

@@ -92,3 +92,4 @@ __SYCL_ASPECT(ext_intel_eus_per_xe_core, 94)
 __SYCL_ASPECT(ext_intel_max_lanes_per_hw_thread, 95)
 __SYCL_ASPECT(ext_oneapi_ipc_physical_memory, 96)
 __SYCL_ASPECT(ext_oneapi_register_host_memory, 97)
+__SYCL_ASPECT(ext_oneapi_per_event_profiling, 98)

sycl/source/CMakeLists.txt

@@ -367,6 +367,7 @@ set(SYCL_COMMON_SOURCES
     "platform.cpp"
     "printers.cpp"
     "queue.cpp"
+    "reusable_events.cpp"
     "sampler.cpp"
     "stream.cpp"
     "virtual_mem.cpp"

sycl/source/detail/context_impl.cpp

@@ -574,6 +574,40 @@ context_impl::get_default_memory_pool(const context &Context,
   return MemPoolImplPtr;
 }
 
+bool context_impl::supportsReusableEvents() {
+  std::lock_guard<std::mutex> Lock{MReusableEventsSupportMutex};
+  if (MReusableEventsSupport) {
+    return *MReusableEventsSupport;
+  }
+
+  MReusableEventsSupport = std::all_of(
+      MDevices.cbegin(), MDevices.cend(), [this](detail::device_impl *DevImpl) {
+        ur_bool_t ReusableEventsSupport = false;
+        ur_result_t Result =
+            getAdapter().call_nocheck<UrApiKind::urDeviceGetInfo>(
+                DevImpl->getHandleRef(),
+                UR_DEVICE_INFO_REUSABLE_EVENTS_SUPPORT_EXP,
+                sizeof(ReusableEventsSupport), &ReusableEventsSupport, nullptr);
+        return (Result == UR_RESULT_SUCCESS) && ReusableEventsSupport;
+      });
+
+  return *MReusableEventsSupport;
+}
+
+bool context_impl::supportsEventProfiling() {
+  std::lock_guard<std::mutex> Lock{MEventProfilingSupportMutex};
+  if (MEventProfilingSupport) {
+    return *MEventProfilingSupport;
+  }
+
+  MEventProfilingSupport = std::all_of(
+      MDevices.cbegin(), MDevices.cend(), [](detail::device_impl *DevImpl) {
+        return DevImpl->has(aspect::ext_oneapi_per_event_profiling);
+      });
+
+  return *MEventProfilingSupport;
+}
+
 } // namespace detail
 } // namespace _V1
 } // namespace sycl

sycl/source/detail/context_impl.hpp

@@ -269,6 +269,9 @@ class context_impl : public std::enable_shared_from_this<context_impl> {
   /// @param UrGraphHandle The native UR graph handle to deregister
   void deregisterNativeGraph(ur_exp_graph_handle_t UrGraphHandle);
 
+  bool supportsReusableEvents();
+  bool supportsEventProfiling();
+
 private:
   bool MOwnedByRuntime;
   async_handler MAsyncHandler;
@@ -278,6 +281,10 @@ class context_impl : public std::enable_shared_from_this<context_impl> {
   property_list MPropList;
   mutable KernelProgramCache MKernelProgramCache;
   mutable PropertySupport MSupportBufferLocationByDevices;
+  std::optional<bool> MReusableEventsSupport;
+  std::mutex MReusableEventsSupportMutex;
+  std::optional<bool> MEventProfilingSupport;
+  std::mutex MEventProfilingSupportMutex;
 
   // Device pools.
   // Weak_ptr preventing circular dependency between memory_pool_impl and

sycl/source/detail/device_impl.hpp

@@ -1637,6 +1637,11 @@ class device_impl {
                  UR_DEVICE_INFO_IPC_PHYSICAL_MEMORY_SUPPORT_EXP>()
           .value_or(0);
     }
+    CASE(ext_oneapi_per_event_profiling) {
+      return get_info_impl_nocheck<
+                 UR_DEVICE_INFO_PER_EVENT_PROFILING_SUPPORT_EXP>()
+          .value_or(0);
+    }
     else {
       return false; // This device aspect has not been implemented yet.
     }

sycl/source/detail/event_impl.cpp

@@ -182,14 +182,70 @@ event_impl::event_impl(HostEventState State, private_tag) : MState(State) {
 
 void event_impl::setQueue(queue_impl &Queue) {
   MQueue = Queue.weak_from_this();
-  MIsProfilingEnabled = Queue.MIsProfilingEnabled;
+  // Inherit profiling from the queue only if not already enabled per-event,
+  // per-event profiling takes precedence over queue-level profiling.
+  MIsProfilingEnabled = MIsProfilingEnabled || Queue.MIsProfilingEnabled;
+}
+
+void event_impl::toDeviceEvent(queue_impl &Queue) {
+  assert(MIsDefaultConstructed);
+
+  initContextIfNeeded();
+  setQueue(Queue);
+
+  ur_event_handle_t EventHandle = nullptr;
+  ur_exp_event_desc_t Desc = {};
+  Desc.stype = UR_STRUCTURE_TYPE_EXP_EVENT_DESC;
+  if (MIsProfilingEnabled) {
+    Desc.flags = UR_EXP_EVENT_FLAG_ENABLE_PROFILING;
+  }
+
+  ur_result_t Result =
+      getAdapter().call_nocheck<sycl::detail::UrApiKind::urEventCreateExp>(
+          MContext->getHandleRef(), Queue.getDeviceImpl().getHandleRef(), &Desc,
+          &EventHandle);
+  if (Result != UR_RESULT_SUCCESS) {
+    throw sycl::exception(sycl::make_error_code(errc::runtime),
+                          "Failed to create an event.");
+  }
 
-  // TODO After setting the queue, the event is no longer default
-  // constructed. Consider a design change which would allow
-  // for such a change regardless of the construction method.
+  setHandle(EventHandle);
   MIsDefaultConstructed = false;
 }
 
+ur_event_handle_t event_impl::getHandleReusable(queue_impl &Queue) {
+  initContextIfNeeded();
+
+  if (MContext->supportsReusableEvents()) {
+    if (MIsDefaultConstructed) {
+      // If the event was constructed (through make_event or a default
+      // constructor), but not enqueued for signaling yet, change the event
+      // state from default constructed to device event.
+      toDeviceEvent(Queue);
+    }
+  } else {
+    // If the context does not support reusable events, then release the
+    // previous event and set the handle to nullptr, so UR can create a new
+    // event during command submission.
+    ur_event_handle_t CurrentHandle = getHandle();
+    if (CurrentHandle != nullptr) {
+      getAdapter().call<UrApiKind::urEventRelease>(CurrentHandle);
+      setHandle(nullptr);
+    }
+  }
+
+  return getHandle();
+}
+
+void event_impl::setHandleReusable(ur_event_handle_t Handle) {
+  // If reusable events are supported do nothing, as the UR handle is already
+  // set. If reusable events are not supported, set the handle of the new UR
+  // event.
+  if (!MContext->supportsReusableEvents()) {
+    setHandle(Handle);
+  }
+}
+
 void event_impl::initHostProfilingInfo() {
   assert(isHost() && "This must be a host event");
   assert(MState == HES_NotComplete &&
@@ -317,7 +373,8 @@ void event_impl::checkProfilingPreconditions() const {
     throw sycl::exception(
         make_error_code(sycl::errc::invalid),
         "Profiling information is unavailable as the queue associated with "
-        "the event does not have the 'enable_profiling' property.");
+        "the event does not have the 'enable_profiling' property or the "
+        "event was not created with the 'enable_profiling' property.");
   }
 }
 

sycl/source/detail/event_impl.hpp

@@ -96,14 +96,28 @@ class event_impl {
                                         private_tag{});
   }
 
-  /// Sets a queue associated with the event
-  ///
-  /// Please note that this function changes the event state
-  /// as it was constructed with the queue based constructor.
+  /// Sets a queue associated with the event.
   ///
   /// \param Queue is a queue to be associated with the event
   void setQueue(queue_impl &Queue);
 
+  /// Converts the event from default constructed to device event.
+  ///
+  /// \param Queue is a queue to be associated with the event
+  void toDeviceEvent(queue_impl &Queue);
+
+  /// Returns an event UR handle and applies additional logic
+  /// related to reusable events.
+  ///
+  /// \param Queue is a queue to be associated with the event
+  ur_event_handle_t getHandleReusable(queue_impl &Queue);
+
+  /// Sets the event UR handle and applies additional logic
+  /// related to reusable events.
+  ///
+  /// \param Handle is a UR handle to be set
+  void setHandleReusable(ur_event_handle_t Handle);
+
   /// Waits for the event.
   ///
   /// \param Success is an optional parameter that, when set to a non-null