[SYCL] Physical memory IPC

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

@@ -33,6 +33,15 @@ __SYCL_EXPORT handle get(void *Ptr, const sycl::context &Ctx);
 __SYCL_EXPORT void put(handle &HandleData, const sycl::context &Ctx);
 } // namespace ext::oneapi::experimental::ipc::memory
 
+namespace ext::oneapi::experimental {
+class physical_mem;
+} // namespace ext::oneapi::experimental
+
+namespace ext::oneapi::experimental::ipc::physical_memory {
+__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 {
 
 using handle_data_t = std::vector<std::byte>;
@@ -59,6 +68,9 @@ struct handle {
   friend __SYCL_EXPORT handle memory::get(void *Ptr, const sycl::context &Ctx);
   friend __SYCL_EXPORT void memory::put(handle &HandleData,
                                         const sycl::context &Ctx);
+  friend __SYCL_EXPORT handle physical_memory::get(physical_mem &physmem);
+  friend __SYCL_EXPORT void physical_memory::put(handle &ipc_handle,
+                                                 const sycl::context &ctx);
 };
 
 } // namespace ext::oneapi::experimental::ipc

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

@@ -0,0 +1,91 @@
+//==------- ipc_memory.hpp -- SYCL inter-process for physical mem ----------==//
+//
+// 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/ext/oneapi/virtual_mem/physical_mem.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 {
+__SYCL_EXPORT sycl::ext::oneapi::experimental::physical_mem
+openIPCPhysicalMemHandle(const std::byte *HandleData, size_t HandleDataSize,
+                         const sycl::context &Ctx, const sycl::device &Dev);
+}
+
+namespace ext::oneapi::experimental::ipc::physical_memory {
+
+__SYCL_EXPORT ipc::handle get(physical_mem &physmem);
+
+__SYCL_EXPORT void put(handle &ipc_handle, const sycl::context &ctx);
+
+inline void put(handle &ipc_handle) {
+  sycl::device Dev;
+  sycl::context Ctx = Dev.get_platform().khr_get_default_context();
+  return put(ipc_handle, Ctx);
+}
+
+inline physical_mem open(const ipc::handle_data_t &HandleData,
+                         const sycl::context &Ctx, const sycl::device &Dev) {
+  return sycl::detail::openIPCPhysicalMemHandle(HandleData.data(),
+                                                HandleData.size(), Ctx, Dev);
+}
+
+inline physical_mem open(ipc::handle_data_t HandleData,
+                         const sycl::device &Dev) {
+  sycl::context Ctx = Dev.get_platform().khr_get_default_context();
+  return open(HandleData, Ctx, Dev);
+}
+
+inline physical_mem open(ipc::handle_data_t HandleData) {
+  sycl::device Dev;
+  sycl::context Ctx = Dev.get_platform().khr_get_default_context();
+  return open(HandleData, Ctx, Dev);
+}
+
+#if __cpp_lib_span
+inline physical_mem open(const ipc::handle_data_view_t &HandleDataView,
+                         const sycl::context &Ctx, const sycl::device &Dev) {
+  return sycl::detail::openIPCPhysicalMemHandle(
+      HandleDataView.data(), HandleDataView.size(), Ctx, Dev);
+}
+
+inline physical_mem open(ipc::handle_data_view_t HandleDataView,
+                         const sycl::device &Dev) {
+  sycl::context Ctx = Dev.get_platform().khr_get_default_context();
+  return open(HandleDataView, Ctx, Dev);
+}
+
+inline physical_mem open(ipc::handle_data_view_t HandleDataView) {
+  sycl::device Dev;
+  sycl::context Ctx = Dev.get_platform().khr_get_default_context();
+  return open(HandleDataView, Ctx, Dev);
+}
+#endif
+
+} // namespace ext::oneapi::experimental::ipc::physical_memory
+} // namespace _V1
+} // namespace sycl
+
+#endif

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

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

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

@@ -26,11 +26,34 @@ 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;
 
 public:
+  template <typename PropertyListT = empty_properties_t>
+  physical_mem(const device &SyclDevice, const context &SyclContext,
+               size_t NumBytes,
+               const PropertyListT &PropList = empty_properties_t{}) {
+
+    bool EnableIPC = PropList.template has_property<enable_ipc_key>();
+
+    create(SyclDevice, SyclContext, NumBytes, EnableIPC);
+  }
+
+  template <typename PropertyListT = empty_properties_t>
+  physical_mem(const queue &SyclQueue, size_t NumBytes,
+               const PropertyListT &PropList = empty_properties_t{})
+      : physical_mem(SyclQueue.get_device(), SyclQueue.get_context(), NumBytes,
+                     PropList) {}
+
   physical_mem(const device &SyclDevice, const context &SyclContext,
                size_t NumBytes);
 
@@ -57,10 +80,19 @@ class __SYCL_EXPORT physical_mem
 
   size_t size() const noexcept;
 
+  bool ipc_enabled() const noexcept;
+
 private:
   std::shared_ptr<sycl::detail::physical_mem_impl> impl;
+  void create(const device &SyclDevice, const context &SyclContext,
+              size_t NumBytes, bool EnableIPC);
+  physical_mem(std::shared_ptr<sycl::detail::physical_mem_impl> Impl)
+      : impl(std::move(Impl)) {}
 };
 
+template <>
+struct is_property_key_of<enable_ipc_key, physical_mem> : std::true_type {};
+
 } // namespace ext::oneapi::experimental
 } // namespace _V1
 } // namespace sycl

sycl/include/sycl/info/aspects.def

@@ -90,3 +90,4 @@ __SYCL_ASPECT(ext_intel_xe_clusters_per_region, 92)
 __SYCL_ASPECT(ext_intel_xe_cores_per_cluster, 93)
 __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/source/CMakeLists.txt

@@ -339,6 +339,7 @@ set(SYCL_COMMON_SOURCES
     "image.cpp"
     "interop_handle.cpp"
     "ipc_memory.cpp"
+    "ipc_physical_memory.cpp"
     "kernel.cpp"
     "kernel_bundle.cpp"
     "physical_mem.cpp"

sycl/source/detail/device_impl.hpp

@@ -1612,6 +1612,14 @@ class device_impl {
       return get_info_impl_nocheck<UR_DEVICE_INFO_DEVICE_WAIT_SUPPORT_EXP>()
           .value_or(0);
     }
+    // TODO IPC physical memory
+    /*
+    CASE(ext_oneapi_ipc_physical_memory) {
+      return
+    get_info_impl_nocheck<UR_DEVICE_INFO_IPC_PHYSICAL_MEMORY_SUPPORT_EXP>()
+          .value_or(0);
+    }
+    */
     else {
       return false; // This device aspect has not been implemented yet.
     }

sycl/source/detail/physical_mem_impl.hpp

@@ -38,11 +38,18 @@ inline ur_virtual_mem_access_flag_t AccessModeToVirtualAccessFlags(
 class physical_mem_impl {
 public:
   physical_mem_impl(device_impl &DeviceImpl, const context &SyclContext,
-                    size_t NumBytes)
+                    size_t NumBytes, bool EnableIPC = false)
       : MDevice(DeviceImpl), MContext(getSyclObjImpl(SyclContext)),
-        MNumBytes(NumBytes) {
+        MNumBytes(NumBytes), MIPCEnabled(EnableIPC) {
     adapter_impl &Adapter = MContext->getAdapter();
 
+    // TODO IPC physical memory
+    // TODO: Pass IPC flags to UR once the API is available
+    // ur_physical_mem_properties_t Props{};
+    // if (EnableIPC) {
+    //   Props.flags = UR_PHYSICAL_MEM_FLAG_IPC_EXPORTABLE;
+    // }
+
     auto Err = Adapter.call_nocheck<UrApiKind::urPhysicalMemCreate>(
         MContext->getHandleRef(), MDevice.getHandleRef(), MNumBytes, nullptr,
         &MPhysicalMem);
@@ -55,9 +62,22 @@ class physical_mem_impl {
     Adapter.checkUrResult(Err);
   }
 
+  physical_mem_impl(device_impl &DeviceImpl, const context &SyclContext,
+                    size_t NumBytes, ur_physical_mem_handle_t PhysicalMem)
+      : MPhysicalMem(PhysicalMem), MDevice(DeviceImpl),
+        MContext(getSyclObjImpl(SyclContext)), MNumBytes(NumBytes),
+        MOpenedFromIpc(true) {}
+
   ~physical_mem_impl() noexcept(false) {
     adapter_impl &Adapter = MContext->getAdapter();
-    Adapter.call<UrApiKind::urPhysicalMemRelease>(MPhysicalMem);
+
+    if (MOpenedFromIpc) {
+      // TODO IPC physical memory
+      // Adapter.call<UrApiKind::urIPCClosePhysMemHandleExp>(
+      //    MContext->getHandleRef(), MPhysicalMem);
+    } else {
+      Adapter.call<UrApiKind::urPhysicalMemRelease>(MPhysicalMem);
+    }
   }
 
   void *map(uintptr_t Ptr, size_t NumBytes,
@@ -77,6 +97,7 @@ class physical_mem_impl {
   }
   device get_device() const { return createSyclObjFromImpl<device>(MDevice); }
   size_t size() const noexcept { return MNumBytes; }
+  bool ipc_enabled() const noexcept { return MIPCEnabled; }
 
   ur_physical_mem_handle_t &getHandleRef() { return MPhysicalMem; }
   const ur_physical_mem_handle_t &getHandleRef() const { return MPhysicalMem; }
@@ -86,6 +107,8 @@ class physical_mem_impl {
   device_impl &MDevice;
   const std::shared_ptr<context_impl> MContext;
   const size_t MNumBytes;
+  const bool MIPCEnabled = false;
+  const bool MOpenedFromIpc = false;
 };
 
 } // namespace detail

sycl/source/ipc_physical_memory.cpp

@@ -0,0 +1,120 @@
+//==------- ipc_physical_memory.cpp -- SYCL inter-process for physical mem -==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#include <detail/adapter_impl.hpp>
+#include <detail/context_impl.hpp>
+#include <detail/physical_mem_impl.hpp>
+#include <sycl/context.hpp>
+#include <sycl/ext/oneapi/experimental/ipc_physical_memory.hpp>
+
+namespace sycl {
+inline namespace _V1 {
+
+namespace detail {
+
+__SYCL_EXPORT sycl::ext::oneapi::experimental::physical_mem
+openIPCPhysicalMemHandle(const std::byte *HandleData, size_t HandleDataSize,
+                         const sycl::context &Ctx, const sycl::device &Dev) {
+  if (!Dev.has(aspect::ext_oneapi_ipc_physical_memory))
+    throw sycl::exception(
+        sycl::make_error_code(errc::feature_not_supported),
+        "Device does not support aspect::ext_oneapi_ipc_physical_memory.");
+
+  auto CtxImpl = sycl::detail::getSyclObjImpl(Ctx);
+  sycl::detail::adapter_impl &Adapter = CtxImpl->getAdapter();
+
+  ur_physical_mem_handle_t PhysMemHandle = nullptr;
+  ur_result_t UrRes = UR_RESULT_SUCCESS;
+  // TODO IPC physical memory
+  /*
+  UrRes =
+  Adapter.call_nocheck<sycl::detail::UrApiKind::urIPCOpenPhysMemHandleExp>(
+      CtxImpl->getHandleRef(), getSyclObjImpl(Dev)->getHandleRef(),
+      HandleData, HandleDataSize, &PhysMemHandle);
+  */
+  if (UrRes == UR_RESULT_ERROR_INVALID_VALUE)
+    throw sycl::exception(
+        sycl::make_error_code(errc::invalid),
+        "HandleData data size does not correspond to the target platform's "
+        "IPC physical memory handle size.");
+  Adapter.checkUrResult(UrRes);
+  if (PhysMemHandle == nullptr)
+    throw sycl::exception(
+        sycl::make_error_code(errc::runtime),
+        "urIPCOpenPhysMemHandleExp returned success but did not produce a "
+        "valid physical memory handle.");
+
+  try {
+    // Query the actual allocation size from the opened handle so that
+    // physical_mem::size() returns the correct value.
+    size_t NumBytes = 0;
+    Adapter.call<sycl::detail::UrApiKind::urPhysicalMemGetInfo>(
+        PhysMemHandle, UR_PHYSICAL_MEM_INFO_SIZE, sizeof(size_t), &NumBytes,
+        nullptr);
+
+    auto PhysMemImpl = std::make_shared<sycl::detail::physical_mem_impl>(
+        *getSyclObjImpl(Dev), Ctx, NumBytes, PhysMemHandle);
+    return sycl::detail::createSyclObjFromImpl<
+        ext::oneapi::experimental::physical_mem>(PhysMemImpl);
+  } catch (...) {
+    // TODO IPC physical memory
+    /*
+    Adapter.call_nocheck<sycl::detail::UrApiKind::urIPCClosePhysMemHandleExp>(
+        CtxImpl->getHandleRef(), PhysMemHandle);
+    */
+    throw;
+  }
+}
+
+} // namespace detail
+
+namespace ext::oneapi::experimental::ipc::physical_memory {
+
+__SYCL_EXPORT handle get(physical_mem &physmem) {
+  if (!physmem.ipc_enabled())
+    throw sycl::exception(
+        sycl::make_error_code(errc::invalid),
+        "physical_mem was not created with inter-process sharing enabled "
+        "via the enable_ipc property.");
+
+  auto PhysMemImpl = sycl::detail::getSyclObjImpl(physmem);
+  auto CtxImpl = sycl::detail::getSyclObjImpl(physmem.get_context());
+  sycl::detail::adapter_impl &Adapter = CtxImpl->getAdapter();
+
+  void *HandlePtr = nullptr;
+  size_t HandleSize = 0;
+  auto UrRes = UR_RESULT_SUCCESS;
+  // TODO IPC physical memory
+  // UR returns total handle size (including the extension)
+  /*
+  UrRes =
+  Adapter.call_nocheck<sycl::detail::UrApiKind::urIPCGetPhysMemHandleExp>(
+    CtxImpl->getHandleRef(), PhysMemImpl->getHandleRef(), &HandlePtr,
+    &HandleSize);
+  */
+
+  return {HandlePtr, HandleSize};
+}
+
+__SYCL_EXPORT void put(handle &ipc_handle, const sycl::context &ctx) {
+  auto CtxImpl = sycl::detail::getSyclObjImpl(ctx);
+  sycl::detail::adapter_impl &Adapter = CtxImpl->getAdapter();
+
+  ur_result_t UrRes = UR_RESULT_SUCCESS;
+  // TODO IPC physical memory
+  /*
+  UrRes =
+  Adapter.call_nocheck<sycl::detail::UrApiKind::urIPCPutPhysMemHandleExp>(
+      CtxImpl->getHandleRef(), ipc_handle.MData);
+  */
+  Adapter.checkUrResult(UrRes);
+}
+
+} // namespace ext::oneapi::experimental::ipc::physical_memory
+} // namespace _V1
+} // namespace sycl