From 3c06e6d4a006adb934db0a8ea367850149157fee Mon Sep 17 00:00:00 2001 From: "Ptak, Slawomir" Date: Fri, 15 May 2026 12:13:29 +0000 Subject: [PATCH 1/6] Initial commit. --- .../oneapi/experimental/detail/ipc_common.hpp | 9 ++ .../experimental/ipc_physical_memory.hpp | 91 +++++++++++++++++++ .../sycl/ext/oneapi/properties/property.hpp | 1 + .../ext/oneapi/virtual_mem/physical_mem.hpp | 28 +++++- sycl/include/sycl/info/aspects.def | 1 + sycl/source/CMakeLists.txt | 1 + sycl/source/detail/device_impl.hpp | 8 ++ sycl/source/detail/physical_mem_impl.hpp | 13 ++- sycl/source/ipc_physical_memory.cpp | 45 +++++++++ sycl/source/physical_mem.cpp | 28 +++--- 10 files changed, 208 insertions(+), 17 deletions(-) create mode 100644 sycl/include/sycl/ext/oneapi/experimental/ipc_physical_memory.hpp create mode 100644 sycl/source/ipc_physical_memory.cpp diff --git a/sycl/include/sycl/ext/oneapi/experimental/detail/ipc_common.hpp b/sycl/include/sycl/ext/oneapi/experimental/detail/ipc_common.hpp index a6cf6d481b8c1..42433c7f7e13e 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/detail/ipc_common.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/detail/ipc_common.hpp @@ -33,6 +33,14 @@ __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); +} // namespace ext::oneapi::experimental::ipc::physical_memory + namespace ext::oneapi::experimental::ipc { using handle_data_t = std::vector; @@ -59,6 +67,7 @@ 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); }; } // namespace ext::oneapi::experimental::ipc diff --git a/sycl/include/sycl/ext/oneapi/experimental/ipc_physical_memory.hpp b/sycl/include/sycl/ext/oneapi/experimental/ipc_physical_memory.hpp new file mode 100644 index 0000000000000..96aa151a0b75e --- /dev/null +++ b/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 +#include +#include +#include +#include +#include + +#include "detail/ipc_common.hpp" + +#include + +#if __has_include() +#include +#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 diff --git a/sycl/include/sycl/ext/oneapi/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp index d73fdd6e51666..409086c48f76c 100644 --- a/sycl/include/sycl/ext/oneapi/properties/property.hpp +++ b/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 struct PropertyToKind { diff --git a/sycl/include/sycl/ext/oneapi/virtual_mem/physical_mem.hpp b/sycl/include/sycl/ext/oneapi/virtual_mem/physical_mem.hpp index aecef93473911..732442596fff2 100644 --- a/sycl/include/sycl/ext/oneapi/virtual_mem/physical_mem.hpp +++ b/sycl/include/sycl/ext/oneapi/virtual_mem/physical_mem.hpp @@ -26,17 +26,33 @@ 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; +}; + +inline constexpr enable_ipc_key::value_t enable_ipc; + class __SYCL_EXPORT physical_mem : public sycl::detail::OwnerLessBase { friend sycl::detail::ImplUtils; public: + template physical_mem(const device &SyclDevice, const context &SyclContext, - size_t NumBytes); + size_t NumBytes, + const PropertyListT &PropList = empty_properties_t{}) { - physical_mem(const queue &SyclQueue, size_t NumBytes) - : physical_mem(SyclQueue.get_device(), SyclQueue.get_context(), - NumBytes) {} + bool EnableIPC = PropertyListT::template has_property(); + + create(SyclDevice, SyclContext, NumBytes, EnableIPC); + } + + template + 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 physical_mem &rhs) = default; physical_mem(physical_mem &&rhs) = default; @@ -57,8 +73,12 @@ class __SYCL_EXPORT physical_mem size_t size() const noexcept; + bool ipc_enabled() const; + private: std::shared_ptr impl; + void create(const device &SyclDevice, const context &SyclContext, + size_t NumBytes, bool EnableIPC); }; } // namespace ext::oneapi::experimental diff --git a/sycl/include/sycl/info/aspects.def b/sycl/include/sycl/info/aspects.def index 78785ccf8d647..69ff4e02ee761 100644 --- a/sycl/include/sycl/info/aspects.def +++ b/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) diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt index d197381a1a307..6e4d857f1c59b 100644 --- a/sycl/source/CMakeLists.txt +++ b/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" diff --git a/sycl/source/detail/device_impl.hpp b/sycl/source/detail/device_impl.hpp index f2d461c007ea1..f334f41fd60be 100644 --- a/sycl/source/detail/device_impl.hpp +++ b/sycl/source/detail/device_impl.hpp @@ -1612,6 +1612,14 @@ class device_impl { return get_info_impl_nocheck() .value_or(0); } + // TODO IPC physical memory + /* + CASE(ext_oneapi_ipc_physical_memory) { + return + get_info_impl_nocheck() + .value_or(0); + } + */ else { return false; // This device aspect has not been implemented yet. } diff --git a/sycl/source/detail/physical_mem_impl.hpp b/sycl/source/detail/physical_mem_impl.hpp index 3bc4d6865187c..0da1032c43dbc 100644 --- a/sycl/source/detail/physical_mem_impl.hpp +++ b/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( MContext->getHandleRef(), MDevice.getHandleRef(), MNumBytes, nullptr, &MPhysicalMem); @@ -77,6 +84,7 @@ class physical_mem_impl { } device get_device() const { return createSyclObjFromImpl(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 +94,7 @@ class physical_mem_impl { device_impl &MDevice; const std::shared_ptr MContext; const size_t MNumBytes; + const bool MIPCEnabled; }; } // namespace detail diff --git a/sycl/source/ipc_physical_memory.cpp b/sycl/source/ipc_physical_memory.cpp new file mode 100644 index 0000000000000..9d16f12dcde25 --- /dev/null +++ b/sycl/source/ipc_physical_memory.cpp @@ -0,0 +1,45 @@ +//==------- 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 +#include +#include +#include + +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."); + + return sycl::ext::oneapi::experimental::physical_mem{Dev, Ctx, 0}; +} + +} // namespace detail + +namespace ext::oneapi::experimental::ipc::physical_memory { + +__SYCL_EXPORT handle get(physical_mem &physmem) { + void *HandlePtr = nullptr; + size_t HandleSize = 0; + + return {HandlePtr, HandleSize}; +} + +__SYCL_EXPORT void put(handle &ipc_handle, const sycl::context &ctx) {} + +} // namespace ext::oneapi::experimental::ipc::physical_memory +} // namespace _V1 +} // namespace sycl diff --git a/sycl/source/physical_mem.cpp b/sycl/source/physical_mem.cpp index 67486c83df317..ca7729f21135e 100644 --- a/sycl/source/physical_mem.cpp +++ b/sycl/source/physical_mem.cpp @@ -13,17 +13,6 @@ namespace sycl { inline namespace _V1 { namespace ext::oneapi::experimental { -physical_mem::physical_mem(const device &SyclDevice, const context &SyclContext, - size_t NumBytes) { - if (!SyclDevice.has(aspect::ext_oneapi_virtual_mem)) - throw sycl::exception( - sycl::make_error_code(sycl::errc::feature_not_supported), - "Device does not support aspect::ext_oneapi_virtual_mem."); - - impl = std::make_shared( - *detail::getSyclObjImpl(SyclDevice), SyclContext, NumBytes); -} - void *physical_mem::map(uintptr_t Ptr, size_t NumBytes, address_access_mode Mode, size_t Offset) const { return impl->map(Ptr, NumBytes, Mode, Offset); @@ -33,6 +22,23 @@ context physical_mem::get_context() const { return impl->get_context(); } device physical_mem::get_device() const { return impl->get_device(); } size_t physical_mem::size() const noexcept { return impl->size(); } +void physical_mem::create(const device &SyclDevice, const context &SyclContext, + size_t NumBytes, bool EnableIPC) { + if (!SyclDevice.has(aspect::ext_oneapi_virtual_mem)) + throw sycl::exception( + sycl::make_error_code(sycl::errc::feature_not_supported), + "Device does not support aspect::ext_oneapi_virtual_mem."); + + if (EnableIPC && !SyclDevice.has(aspect::ext_oneapi_ipc_physical_memory)) { + throw sycl::exception( + sycl::make_error_code(sycl::errc::feature_not_supported), + "Device does not support aspect::ext_oneapi_ipc_physical_memory."); + } + + impl = std::make_shared( + *detail::getSyclObjImpl(SyclDevice), SyclContext, NumBytes, EnableIPC); +} + } // namespace ext::oneapi::experimental } // namespace _V1 } // namespace sycl From eecbfcf482def3a9f0d315e35f856ea8f290ce9d Mon Sep 17 00:00:00 2001 From: "Ptak, Slawomir" Date: Mon, 25 May 2026 08:55:39 +0000 Subject: [PATCH 2/6] Draft of IPC for physical memory functions, minor updates. --- .../oneapi/experimental/detail/ipc_common.hpp | 3 + .../ext/oneapi/virtual_mem/physical_mem.hpp | 9 ++- sycl/source/detail/physical_mem_impl.hpp | 18 ++++- sycl/source/ipc_physical_memory.cpp | 76 ++++++++++++++++++- sycl/source/physical_mem.cpp | 1 + 5 files changed, 101 insertions(+), 6 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/detail/ipc_common.hpp b/sycl/include/sycl/ext/oneapi/experimental/detail/ipc_common.hpp index 42433c7f7e13e..a16d1919b6b83 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/detail/ipc_common.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/detail/ipc_common.hpp @@ -39,6 +39,7 @@ class physical_mem; 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 { @@ -68,6 +69,8 @@ struct handle { 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 diff --git a/sycl/include/sycl/ext/oneapi/virtual_mem/physical_mem.hpp b/sycl/include/sycl/ext/oneapi/virtual_mem/physical_mem.hpp index 732442596fff2..935d01286fde7 100644 --- a/sycl/include/sycl/ext/oneapi/virtual_mem/physical_mem.hpp +++ b/sycl/include/sycl/ext/oneapi/virtual_mem/physical_mem.hpp @@ -73,14 +73,19 @@ class __SYCL_EXPORT physical_mem size_t size() const noexcept; - bool ipc_enabled() const; + bool ipc_enabled() const noexcept; private: std::shared_ptr impl; void create(const device &SyclDevice, const context &SyclContext, - size_t NumBytes, bool EnableIPC); + size_t NumBytes, bool EnableIPC); + physical_mem(std::shared_ptr Impl) + : impl(std::move(Impl)) {} }; +template <> +struct is_property_key_of : std::true_type {}; + } // namespace ext::oneapi::experimental } // namespace _V1 } // namespace sycl diff --git a/sycl/source/detail/physical_mem_impl.hpp b/sycl/source/detail/physical_mem_impl.hpp index 0da1032c43dbc..a7677b09b1ed3 100644 --- a/sycl/source/detail/physical_mem_impl.hpp +++ b/sycl/source/detail/physical_mem_impl.hpp @@ -62,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(MPhysicalMem); + + if (MOpenedFromIpc) { + // TODO IPC physical memory + // Adapter.call( + // MContext->getHandleRef(), MPhysicalMem); + } else { + Adapter.call(MPhysicalMem); + } } void *map(uintptr_t Ptr, size_t NumBytes, @@ -94,7 +107,8 @@ class physical_mem_impl { device_impl &MDevice; const std::shared_ptr MContext; const size_t MNumBytes; - const bool MIPCEnabled; + const bool MIPCEnabled = false; + const bool MOpenedFromIpc = false; }; } // namespace detail diff --git a/sycl/source/ipc_physical_memory.cpp b/sycl/source/ipc_physical_memory.cpp index 9d16f12dcde25..547b8ad654311 100644 --- a/sycl/source/ipc_physical_memory.cpp +++ b/sycl/source/ipc_physical_memory.cpp @@ -8,6 +8,7 @@ #include #include +#include #include #include @@ -24,7 +25,40 @@ openIPCPhysicalMemHandle(const std::byte *HandleData, size_t HandleDataSize, sycl::make_error_code(errc::feature_not_supported), "Device does not support aspect::ext_oneapi_ipc_physical_memory."); - return sycl::ext::oneapi::experimental::physical_mem{Dev, Ctx, 0}; + auto CtxImpl = sycl::detail::getSyclObjImpl(Ctx); + sycl::detail::adapter_impl &Adapter = CtxImpl->getAdapter(); + + size_t NumBytes; + std::memcpy(&NumBytes, HandleData, sizeof(size_t)); + + const std::byte *ActualHandleData = HandleData + sizeof(size_t); + size_t ActualHandleSize = HandleDataSize - sizeof(size_t); + + ur_physical_mem_handle_t PhysMemHandle = nullptr; + ur_result_t UrRes = UR_RESULT_SUCCESS; + // TODO IPC physical memory + /* + UrRes = + Adapter.call_nocheck( + CtxImpl->getHandleRef(), getSyclObjImpl(Dev)->getHandleRef(), + ActualHandleData, ActualHandleSize, &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."); + + auto PhysMemImpl = std::make_shared( + *getSyclObjImpl(Dev), Ctx, NumBytes, PhysMemHandle); + return sycl::detail::createSyclObjFromImpl< + ext::oneapi::experimental::physical_mem>(PhysMemImpl); } } // namespace detail @@ -32,13 +66,51 @@ openIPCPhysicalMemHandle(const std::byte *HandleData, size_t HandleDataSize, 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; + // Additional memory space required to store the physical memory size. This + // space is allocated by UR at the beginning of the whole handle allocation. + size_t HandleExtensionSize = sizeof(size_t); + auto UrRes = UR_RESULT_SUCCESS; + // TODO IPC physical memory + // UR returns total handle size (including the extension) + /* + UrRes = + Adapter.call_nocheck( + CtxImpl->getHandleRef(), PhysMemImpl->getHandleRef(), &HandlePtr, + &HandleSize, HandleExtensionSize); + */ + + // Copy the physical memory allocation size to the handle + size_t NumBytes = physmem.size(); + std::memcpy(HandlePtr, &NumBytes, sizeof(size_t)); return {HandlePtr, HandleSize}; } -__SYCL_EXPORT void put(handle &ipc_handle, const sycl::context &ctx) {} +__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( + CtxImpl->getHandleRef(), ipc_handle.MData); + */ + Adapter.checkUrResult(UrRes); +} } // namespace ext::oneapi::experimental::ipc::physical_memory } // namespace _V1 diff --git a/sycl/source/physical_mem.cpp b/sycl/source/physical_mem.cpp index ca7729f21135e..7ac2b4ec443bf 100644 --- a/sycl/source/physical_mem.cpp +++ b/sycl/source/physical_mem.cpp @@ -21,6 +21,7 @@ void *physical_mem::map(uintptr_t Ptr, size_t NumBytes, context physical_mem::get_context() const { return impl->get_context(); } device physical_mem::get_device() const { return impl->get_device(); } size_t physical_mem::size() const noexcept { return impl->size(); } +bool physical_mem::ipc_enabled() const noexcept { return impl->ipc_enabled(); } void physical_mem::create(const device &SyclDevice, const context &SyclContext, size_t NumBytes, bool EnableIPC) { From 41d202cc19459703f48064e45e9deb9d5a4d3abd Mon Sep 17 00:00:00 2001 From: "Ptak, Slawomir" Date: Tue, 26 May 2026 09:43:51 +0000 Subject: [PATCH 3/6] Get the physical memory allocation size using UR API when opening the IPC handle. --- sycl/source/ipc_physical_memory.cpp | 37 ++++++++++++++++------------- 1 file changed, 20 insertions(+), 17 deletions(-) diff --git a/sycl/source/ipc_physical_memory.cpp b/sycl/source/ipc_physical_memory.cpp index 547b8ad654311..ca94f59757f78 100644 --- a/sycl/source/ipc_physical_memory.cpp +++ b/sycl/source/ipc_physical_memory.cpp @@ -28,12 +28,6 @@ openIPCPhysicalMemHandle(const std::byte *HandleData, size_t HandleDataSize, auto CtxImpl = sycl::detail::getSyclObjImpl(Ctx); sycl::detail::adapter_impl &Adapter = CtxImpl->getAdapter(); - size_t NumBytes; - std::memcpy(&NumBytes, HandleData, sizeof(size_t)); - - const std::byte *ActualHandleData = HandleData + sizeof(size_t); - size_t ActualHandleSize = HandleDataSize - sizeof(size_t); - ur_physical_mem_handle_t PhysMemHandle = nullptr; ur_result_t UrRes = UR_RESULT_SUCCESS; // TODO IPC physical memory @@ -41,7 +35,7 @@ openIPCPhysicalMemHandle(const std::byte *HandleData, size_t HandleDataSize, UrRes = Adapter.call_nocheck( CtxImpl->getHandleRef(), getSyclObjImpl(Dev)->getHandleRef(), - ActualHandleData, ActualHandleSize, &PhysMemHandle); + HandleData, HandleDataSize, &PhysMemHandle); */ if (UrRes == UR_RESULT_ERROR_INVALID_VALUE) throw sycl::exception( @@ -55,10 +49,26 @@ openIPCPhysicalMemHandle(const std::byte *HandleData, size_t HandleDataSize, "urIPCOpenPhysMemHandleExp returned success but did not produce a " "valid physical memory handle."); - auto PhysMemImpl = std::make_shared( + 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( + PhysMemHandle, UR_PHYSICAL_MEM_INFO_SIZE, sizeof(size_t), &NumBytes, + nullptr); + + auto PhysMemImpl = std::make_shared( *getSyclObjImpl(Dev), Ctx, NumBytes, PhysMemHandle); - return sycl::detail::createSyclObjFromImpl< + return sycl::detail::createSyclObjFromImpl< ext::oneapi::experimental::physical_mem>(PhysMemImpl); + } catch (...) { + // TODO IPC physical memory + /* + Adapter.call_nocheck( + CtxImpl->getHandleRef(), PhysMemHandle); + */ + throw; + } } } // namespace detail @@ -78,9 +88,6 @@ __SYCL_EXPORT handle get(physical_mem &physmem) { void *HandlePtr = nullptr; size_t HandleSize = 0; - // Additional memory space required to store the physical memory size. This - // space is allocated by UR at the beginning of the whole handle allocation. - size_t HandleExtensionSize = sizeof(size_t); auto UrRes = UR_RESULT_SUCCESS; // TODO IPC physical memory // UR returns total handle size (including the extension) @@ -88,13 +95,9 @@ __SYCL_EXPORT handle get(physical_mem &physmem) { UrRes = Adapter.call_nocheck( CtxImpl->getHandleRef(), PhysMemImpl->getHandleRef(), &HandlePtr, - &HandleSize, HandleExtensionSize); + &HandleSize); */ - // Copy the physical memory allocation size to the handle - size_t NumBytes = physmem.size(); - std::memcpy(HandlePtr, &NumBytes, sizeof(size_t)); - return {HandlePtr, HandleSize}; } From 13a9ffeb0ce2124d462e56d2c0fd6d09ed2c1ff7 Mon Sep 17 00:00:00 2001 From: "Ptak, Slawomir" Date: Tue, 26 May 2026 09:45:07 +0000 Subject: [PATCH 4/6] Fix formatting. --- sycl/source/ipc_physical_memory.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/source/ipc_physical_memory.cpp b/sycl/source/ipc_physical_memory.cpp index ca94f59757f78..13eace68eec63 100644 --- a/sycl/source/ipc_physical_memory.cpp +++ b/sycl/source/ipc_physical_memory.cpp @@ -58,9 +58,9 @@ openIPCPhysicalMemHandle(const std::byte *HandleData, size_t HandleDataSize, nullptr); auto PhysMemImpl = std::make_shared( - *getSyclObjImpl(Dev), Ctx, NumBytes, PhysMemHandle); + *getSyclObjImpl(Dev), Ctx, NumBytes, PhysMemHandle); return sycl::detail::createSyclObjFromImpl< - ext::oneapi::experimental::physical_mem>(PhysMemImpl); + ext::oneapi::experimental::physical_mem>(PhysMemImpl); } catch (...) { // TODO IPC physical memory /* From b56db71f8d8527cfcb045bbddb67532f616f3ffe Mon Sep 17 00:00:00 2001 From: "Ptak, Slawomir" Date: Tue, 26 May 2026 10:50:32 +0000 Subject: [PATCH 5/6] Fix unused parameter. --- sycl/include/sycl/ext/oneapi/virtual_mem/physical_mem.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/sycl/ext/oneapi/virtual_mem/physical_mem.hpp b/sycl/include/sycl/ext/oneapi/virtual_mem/physical_mem.hpp index 935d01286fde7..96c2762516b84 100644 --- a/sycl/include/sycl/ext/oneapi/virtual_mem/physical_mem.hpp +++ b/sycl/include/sycl/ext/oneapi/virtual_mem/physical_mem.hpp @@ -43,7 +43,7 @@ class __SYCL_EXPORT physical_mem size_t NumBytes, const PropertyListT &PropList = empty_properties_t{}) { - bool EnableIPC = PropertyListT::template has_property(); + bool EnableIPC = PropList.template has_property(); create(SyclDevice, SyclContext, NumBytes, EnableIPC); } From 0b5fadf74abdedba449db5d6ced5effc915472b4 Mon Sep 17 00:00:00 2001 From: "Ptak, Slawomir" Date: Wed, 27 May 2026 11:18:29 +0000 Subject: [PATCH 6/6] Add non-templated physical_mem constructors. --- sycl/include/sycl/ext/oneapi/virtual_mem/physical_mem.hpp | 7 +++++++ sycl/source/physical_mem.cpp | 5 +++++ 2 files changed, 12 insertions(+) diff --git a/sycl/include/sycl/ext/oneapi/virtual_mem/physical_mem.hpp b/sycl/include/sycl/ext/oneapi/virtual_mem/physical_mem.hpp index 96c2762516b84..3af65f44e8207 100644 --- a/sycl/include/sycl/ext/oneapi/virtual_mem/physical_mem.hpp +++ b/sycl/include/sycl/ext/oneapi/virtual_mem/physical_mem.hpp @@ -54,6 +54,13 @@ class __SYCL_EXPORT physical_mem : physical_mem(SyclQueue.get_device(), SyclQueue.get_context(), NumBytes, PropList) {} + physical_mem(const device &SyclDevice, const context &SyclContext, + size_t NumBytes); + + physical_mem(const queue &SyclQueue, size_t NumBytes) + : physical_mem(SyclQueue.get_device(), SyclQueue.get_context(), + NumBytes) {} + physical_mem(const physical_mem &rhs) = default; physical_mem(physical_mem &&rhs) = default; diff --git a/sycl/source/physical_mem.cpp b/sycl/source/physical_mem.cpp index 7ac2b4ec443bf..408d9f36233ec 100644 --- a/sycl/source/physical_mem.cpp +++ b/sycl/source/physical_mem.cpp @@ -13,6 +13,11 @@ namespace sycl { inline namespace _V1 { namespace ext::oneapi::experimental { +physical_mem::physical_mem(const device &SyclDevice, const context &SyclContext, + size_t NumBytes) { + create(SyclDevice, SyclContext, NumBytes, false); +} + void *physical_mem::map(uintptr_t Ptr, size_t NumBytes, address_access_mode Mode, size_t Offset) const { return impl->map(Ptr, NumBytes, Mode, Offset);