[SYCL] Implement IPC event sharing

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_ipc_event : Aspect<"ext_oneapi_ipc_event">;
 
 // 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_ipc_event],
     []>;
 // This definition serves the only purpose of testing whether the deprecated aspect list defined in here and in SYCL RT
 // match.

sycl/include/sycl/event.hpp

@@ -147,6 +147,10 @@ class __SYCL_EXPORT event : public detail::OwnerLessBase<event> {
   /// \return the backend associated with this platform
   backend get_backend() const noexcept;
 
+  /// Returns true if the event was created with inter-process sharing
+  /// enabled.
+  bool ext_oneapi_ipc_enabled() const noexcept;
+
 private:
   event(std::shared_ptr<detail::event_impl> EventImpl);
 

sycl/include/sycl/ext/oneapi/experimental/detail/ipc_common.hpp

@@ -8,11 +8,33 @@
 
 #pragma once
 
+#include <sycl/detail/defines_elementary.hpp>
+#include <sycl/ext/oneapi/properties/properties.hpp>
+
+namespace sycl {
+inline namespace _V1 {
+
+namespace ext::oneapi::experimental {
+
+// Shared by the inter-process-communication extensions (physical_mem and
+// event); each target adds an is_property_key_of specialisation in its own
+// header.
+struct enable_ipc_key
+    : detail::compile_time_property_key<detail::PropKind::EnableIPC> {
+  using value_t = property_value<enable_ipc_key>;
+};
+
+inline constexpr enable_ipc_key::value_t enable_ipc;
+
+} // namespace ext::oneapi::experimental
+} // namespace _V1
+} // namespace sycl
+
 #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
 
 #include <sycl/context.hpp>
-#include <sycl/detail/defines_elementary.hpp>
 #include <sycl/detail/export.hpp>
+#include <sycl/event.hpp>
 
 #include <cstddef>
 #include <vector>
@@ -42,6 +64,11 @@ __SYCL_EXPORT handle get(physical_mem &physmem);
 __SYCL_EXPORT void put(handle &ipc_handle, const sycl::context &ctx);
 } // namespace ext::oneapi::experimental::ipc::physical_memory
 
+namespace ext::oneapi::experimental::ipc::event {
+__SYCL_EXPORT handle get(const sycl::event &Evt);
+__SYCL_EXPORT void put(handle &IpcHandle, const sycl::context &Ctx);
+} // namespace ext::oneapi::experimental::ipc::event
+
 namespace ext::oneapi::experimental::ipc {
 
 using handle_data_t = std::vector<std::byte>;
@@ -71,6 +98,9 @@ struct handle {
   friend __SYCL_EXPORT handle physical_memory::get(physical_mem &physmem);
   friend __SYCL_EXPORT void physical_memory::put(handle &ipc_handle,
                                                  const sycl::context &ctx);
+  friend __SYCL_EXPORT handle event::get(const sycl::event &Evt);
+  friend __SYCL_EXPORT void event::put(handle &IpcHandle,
+                                       const sycl::context &Ctx);
 };
 
 } // namespace ext::oneapi::experimental::ipc

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

@@ -0,0 +1,114 @@
+//==------- ipc_event.hpp -- SYCL inter-process for 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
+
+#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
+
+#include <sycl/context.hpp>
+#include <sycl/detail/defines_elementary.hpp>
+#include <sycl/detail/export.hpp>
+#include <sycl/device.hpp>
+#include <sycl/event.hpp>
+#include <sycl/ext/oneapi/properties/properties.hpp>
+#include <sycl/platform.hpp>
+
+#include "detail/ipc_common.hpp"
+
+#include <cstddef>
+
+#if __has_include(<span>)
+#include <span>
+#endif
+
+namespace sycl {
+inline namespace _V1 {
+
+namespace detail {
+// Bridge from the user-facing ipc::event::open overloads to a single
+// __SYCL_EXPORT'd implementation.
+__SYCL_EXPORT sycl::event openIPCEventHandle(const std::byte *HandleData,
+                                             size_t HandleDataSize,
+                                             const sycl::context &Ctx);
+
+// Producer-side IPC event creation.
+__SYCL_EXPORT sycl::event makeIPCEvent(const sycl::context &Ctx);
+} // namespace detail
+
+namespace ext::oneapi::experimental {
+
+// enable_ipc_key is defined in detail/ipc_common.hpp and shared with
+// physical_mem; tag it as valid on sycl::event too.
+template <>
+struct is_property_key_of<enable_ipc_key, sycl::event> : std::true_type {};
+
+// Minimum-viable make_event overload that handles only the enable_ipc
+// property. The full reusable-events make_event lands with PR #22186; this
+// overload will fold into it when that PR is merged.
+template <typename PropertyListT = empty_properties_t>
+inline sycl::event make_event(const sycl::context &Ctx,
+                              PropertyListT Props = empty_properties_t{}) {
+  static_assert(is_property_list_v<PropertyListT>,
+                "Props must be a sycl::ext::oneapi::experimental::properties");
+  (void)Props;
+
+  if constexpr (PropertyListT::template has_property<enable_ipc_key>()) {
+    return sycl::detail::makeIPCEvent(Ctx);
+  } else {
+    static_assert(
+        sizeof(PropertyListT) == 0,
+        "make_event without the enable_ipc property is not yet supported. "
+        "The full reusable-events make_event lands with PR #22186.");
+    return sycl::event{};
+  }
+}
+
+} // namespace ext::oneapi::experimental
+
+namespace ext::oneapi::experimental::ipc::event {
+
+// get / put are exported from libsycl; declarations live in
+// detail/ipc_common.hpp because they're friended on ipc::handle.
+
+inline void put(handle &IpcHandle) {
+  sycl::device Dev;
+  sycl::context Ctx = Dev.get_platform().khr_get_default_context();
+  return put(IpcHandle, Ctx);
+}
+
+inline sycl::event open(const ipc::handle_data_t &HandleData,
+                        const sycl::context &Ctx) {
+  return sycl::detail::openIPCEventHandle(HandleData.data(), HandleData.size(),
+                                          Ctx);
+}
+
+inline sycl::event open(const ipc::handle_data_t &HandleData) {
+  sycl::device Dev;
+  sycl::context Ctx = Dev.get_platform().khr_get_default_context();
+  return open(HandleData, Ctx);
+}
+
+#if __cpp_lib_span
+inline sycl::event open(const ipc::handle_data_view_t &HandleDataView,
+                        const sycl::context &Ctx) {
+  return sycl::detail::openIPCEventHandle(HandleDataView.data(),
+                                          HandleDataView.size(), Ctx);
+}
+
+inline sycl::event open(const ipc::handle_data_view_t &HandleDataView) {
+  sycl::device Dev;
+  sycl::context Ctx = Dev.get_platform().khr_get_default_context();
+  return open(HandleDataView, Ctx);
+}
+#endif
+
+} // namespace ext::oneapi::experimental::ipc::event
+} // namespace _V1
+} // namespace sycl
+
+#endif

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

@@ -192,7 +192,7 @@ enum PropKind : uint32_t {
   MaximumSize = 47,
   ZeroInit = 48,
   FastLink = 49,
-  PhysicalMemoryEnableIPC = 50,
+  EnableIPC = 50,
   RegisterHostMemoryReadOnly = 51,
   // PropKindSize must always be the last value.
   PropKindSize = 52,

sycl/include/sycl/ext/oneapi/virtual_mem/physical_mem.hpp

@@ -13,6 +13,7 @@
 #include <sycl/detail/common.hpp>
 #include <sycl/detail/owner_less_base.hpp>
 #include <sycl/device.hpp>
+#include <sycl/ext/oneapi/experimental/detail/ipc_common.hpp>
 #include <sycl/queue.hpp>
 
 namespace sycl {
@@ -26,13 +27,6 @@ namespace ext::oneapi::experimental {
 
 enum class address_access_mode : char { none = 0, read = 1, read_write = 2 };
 
-struct enable_ipc_key : detail::compile_time_property_key<
-                            detail::PropKind::PhysicalMemoryEnableIPC> {
-  using value_t = property_value<enable_ipc_key>;
-};
-
-inline constexpr enable_ipc_key::value_t enable_ipc;
-
 class __SYCL_EXPORT physical_mem
     : public sycl::detail::OwnerLessBase<physical_mem> {
   friend sycl::detail::ImplUtils;

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_ipc_event, 98)

sycl/source/CMakeLists.txt

@@ -359,6 +359,7 @@ set(SYCL_COMMON_SOURCES
     "handler.cpp"
     "image.cpp"
     "interop_handle.cpp"
+    "ipc_event.cpp"
     "ipc_memory.cpp"
     "ipc_physical_memory.cpp"
     "kernel.cpp"

sycl/source/detail/device_impl.hpp

@@ -1622,6 +1622,10 @@ class device_impl {
       return get_info_impl_nocheck<UR_DEVICE_INFO_IPC_MEMORY_SUPPORT_EXP>()
           .value_or(0);
     }
+    CASE(ext_oneapi_ipc_event) {
+      return get_info_impl_nocheck<UR_DEVICE_INFO_IPC_EVENT_SUPPORT_EXP>()
+          .value_or(0);
+    }
     CASE(ext_oneapi_device_wait) {
       return get_info_impl_nocheck<UR_DEVICE_INFO_DEVICE_WAIT_SUPPORT_EXP>()
           .value_or(0);

sycl/source/detail/event_impl.cpp

@@ -167,6 +167,16 @@ event_impl::event_impl(ur_event_handle_t Event, const context &SyclContext,
   }
 }
 
+event_impl::event_impl(ur_event_handle_t Event, const context &SyclContext,
+                       IPCRole Role, private_tag tag)
+    : event_impl(Event, SyclContext, tag) {
+  // ext_oneapi_ipc_enabled() is the predicate for "can be exported via
+  // ipc::event::get()"; the spec restricts that to producer-side events
+  // created with the enable_ipc property.
+  MIPCEnabled = (Role == IPCRole::Producer);
+  MOpenedFromIpc = (Role == IPCRole::Imported);
+}
+
 event_impl::event_impl(queue_impl &Queue, private_tag)
     : MQueue{Queue.weak_from_this()},
       MIsProfilingEnabled{Queue.MIsProfilingEnabled} {

sycl/source/detail/event_impl.hpp

@@ -69,6 +69,17 @@ class event_impl {
   /// \param SyclContext is an instance of SYCL context.
   event_impl(ur_event_handle_t Event, const context &SyclContext, private_tag);
 
+  /// IPC role of an event constructed from a UR handle.
+  enum class IPCRole {
+    /// Producer-side event created with enable_ipc.
+    Producer,
+    /// Consumer-side event imported via ipc::event::open().
+    Imported,
+  };
+
+  event_impl(ur_event_handle_t Event, const context &SyclContext, IPCRole Role,
+             private_tag);
+
   event_impl(queue_impl &Queue, private_tag);
   event_impl(HostEventState State, private_tag);
 
@@ -96,6 +107,25 @@ class event_impl {
                                         private_tag{});
   }
 
+  /// Producer-side IPC event factory: adopts a UR handle that backs
+  /// make_event(..., enable_ipc). Reports ext_oneapi_ipc_enabled() == true.
+  static std::shared_ptr<event_impl>
+  create_ipc_producer_event(ur_event_handle_t Event,
+                            const context &SyclContext) {
+    return std::make_shared<event_impl>(Event, SyclContext, IPCRole::Producer,
+                                        private_tag{});
+  }
+
+  /// Consumer-side IPC event factory: adopts a UR handle returned by
+  /// ipc::event::open(). Imported events cannot be re-exported, so they
+  /// report ext_oneapi_ipc_enabled() == false.
+  static std::shared_ptr<event_impl>
+  create_ipc_imported_event(ur_event_handle_t Event,
+                            const context &SyclContext) {
+    return std::make_shared<event_impl>(Event, SyclContext, IPCRole::Imported,
+                                        private_tag{});
+  }
+
   /// Sets a queue associated with the event
   ///
   /// Please note that this function changes the event state
@@ -407,6 +437,25 @@ class event_impl {
   // storage.
   std::vector<std::weak_ptr<event_impl>> MWeakPostCompleteEvents;
 
+  /// Both IPC flags are written exactly once, by the constructor that
+  /// the create_ipc_*_event factories invoke. After the
+  /// shared_ptr<event_impl> is returned to the caller these fields are
+  /// observed only read-only, so concurrent reads from any number of
+  /// threads are safe without synchronisation: the construction-then-
+  /// publish ordering is established by std::make_shared's release on
+  /// the produced shared_ptr.
+
+  /// Backs event::ext_oneapi_ipc_enabled().
+  bool MIPCEnabled = false;
+
+  /// True only for events imported via ipc::event::open().
+  bool MOpenedFromIpc = false;
+
+public:
+  bool isIPCEnabled() const noexcept { return MIPCEnabled; }
+  bool isOpenedFromIpc() const noexcept { return MOpenedFromIpc; }
+
+protected:
   /// Indicates that the task associated with this event has been submitted by
   /// the queue to the device.
   std::atomic<bool> MIsFlushed = false;

sycl/source/detail/ur_device_info_ret_types.inc

@@ -172,6 +172,7 @@ MAP(UR_DEVICE_INFO_MAX_LANES_PER_HW_THREAD, uint32_t)
 // instead.
 MAP(UR_DEVICE_INFO_2D_BLOCK_ARRAY_CAPABILITIES_EXP, ur_exp_device_2d_block_array_capability_flags_t)
 MAP(UR_DEVICE_INFO_IPC_MEMORY_SUPPORT_EXP, ur_bool_t)
+MAP(UR_DEVICE_INFO_IPC_EVENT_SUPPORT_EXP, ur_bool_t)
 MAP(UR_DEVICE_INFO_ASYNC_USM_ALLOCATIONS_SUPPORT_EXP, ur_bool_t)
 MAP(UR_DEVICE_INFO_BINDLESS_IMAGES_1D_USM_SUPPORT_EXP, ur_bool_t)
 MAP(UR_DEVICE_INFO_BINDLESS_IMAGES_2D_USM_SUPPORT_EXP, ur_bool_t)

sycl/source/event.cpp

@@ -114,6 +114,10 @@ backend event::get_backend() const noexcept try {
   std::abort();
 }
 
+bool event::ext_oneapi_ipc_enabled() const noexcept {
+  return impl->isIPCEnabled();
+}
+
 ur_native_handle_t event::getNative() const { return impl->getNative(); }
 
 std::vector<ur_native_handle_t> event::getNativeVector() const {