From 0dd180a8ff8c2c8bfa79b343d6fe1ee463f1c94a Mon Sep 17 00:00:00 2001 From: "Ptak, Slawomir" Date: Thu, 7 May 2026 09:16:10 +0000 Subject: [PATCH 01/17] Refactor the inter process communication extension doc. Add physical memory IPC doc Specify the namespace of handle struct (and aliases). Fix typo. Add a physical_mem properties constructor arg, clarify the device requirement for the open function, add enable_ipc property. Minor review comments. Add const. Minor fixes. Make ipc_memory namespace deprecated and add a new section which puts IPC for memory in ipc::memory namespace. Add a note that the get function throws an exception if phys_mem was not created with the enable_ipc property. Add a note about deprecating the ipc_memory namespace. Address review comments + namespace nesting change with [[deprecated]] Add missing put functions, change the physical memory lifetime description. Describe the ctx argument for put function. Change ext_oneapi_ipc_enabled() to ipc_enabled() in physical_mem. --- ...neapi_inter_process_communication.asciidoc | 755 ++++++++++++++++-- .../sycl_ext_oneapi_virtual_mem.asciidoc | 22 +- 2 files changed, 723 insertions(+), 54 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_inter_process_communication.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_inter_process_communication.asciidoc index 0489ecf0d896..a0823282aac3 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_inter_process_communication.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_inter_process_communication.asciidoc @@ -57,6 +57,10 @@ there to be a per-platform default context even if the core SYCL specification does not provide a convenient way to get it. _{endnote}_] +This extension also depends on the following other SYCL extensions: + +* link:sycl_ext_oneapi_virtual_mem.asciidoc[sycl_ext_oneapi_virtual_mem] +* link:sycl_ext_oneapi_properties.asciidoc[sycl_ext_oneapi_properties] == Status @@ -70,11 +74,11 @@ specification.* == Overview -This extension adds the ability for SYCL programs to share device USM memory -allocations between processes. This is done by the allocating process creating -a new IPC memory handle through the new free functions and transferring the +This extension adds the ability for SYCL programs to share various objects +or USM memory allocations between processes. This is done by creating a new IPC +handle in one process through the new free functions and transferring the returned handle data to the other processes. The other processes can use the -handle data to retrieve the corresponding device USM memory. +handle data to retrieve the corresponding object or allocation. == Specification @@ -99,30 +103,13 @@ implementation supports. feature-test macro always has this value. |=== -=== Extension to `enum class aspect` - -[source] ----- -namespace sycl { -enum class aspect { - ... - ext_oneapi_ipc_memory -} -} ----- - -If a SYCL device has this aspect, that device supports the `get` and `open` -functions specified in the following section. - - -=== Inter-process communicable memory +=== Generic types -This extension adds a new type aliases `handle_data_t` and `handle_data_view_t` -as well as a new IPC memory handle type `handle` under the `ipc_memory` -experimental namespace. +This extension adds the following types which can be used to share SYCL objects +with another process. ``` -namespace sycl::ext::oneapi::experimental::ipc_memory { +namespace sycl::ext::oneapi::experimental::ipc { using handle_data_t = std::vector; @@ -150,8 +137,8 @@ handle_data_t data() const; ---- !==== -_Preconditions:_ The `put` function has not previously been called with this -handle and the USM device memory associated with this handle has not been freed. +_Preconditions:_ The resources associated with this handle have not been +released by the relevant IPC API or by the destruction of the original object. _Returns:_ The handle data associated with the IPC handle object. This data can be transferred to other processes, but cannot be used to recreate a `handle` @@ -165,21 +152,43 @@ handle_data_view_t data_view() const; ---- !==== -_Preconditions:_ The `put` function has not previously been called with this -handle and the USM device memory associated with this handle has not been freed. +_Preconditions:_ The resources associated with this handle have not been +released by the relevant IPC API or by the destruction of the original object. _Returns:_ A `std::span` with a view into the handle data associated with the IPC handle object. This data can be transferred to other processes, but cannot be used to recreate a `handle` object. - |==== -Additionally, this extension adds new free functions under the `ipc_memory` -experimental namespace. +=== Inter-process communicable memory + +This extension allows for the exchange of USM device memory pointers between +processes. The following aspect and set of free functions are specific to USM +memory sharing feature. + +==== Extension to `enum class aspect` for USM memory + +[source] +---- +namespace sycl { +enum class aspect { + ... + ext_oneapi_ipc_memory +} +} +---- + +If a SYCL device has this aspect, that device supports the `get` and `open` +functions specified in the following section. + +==== New functions and types associated with USM memory + +This extension adds new free functions under the `ipc::memory` experimental +namespace. ``` -namespace sycl::ext::oneapi::experimental::ipc_memory { +namespace sycl::ext::oneapi::experimental::ipc::memory { handle get(void *ptr, const sycl::context &ctx); @@ -189,8 +198,8 @@ void put(handle &ipc_handle, const sycl::context &ctx); void put(handle &ipc_handle); -void *open(const handle_data_t &handle_data, const sycl::context &ctx, - const sycl::device &dev); +void *open(const handle_data_t &handle_data, + const sycl::context &ctx, const sycl::device &dev); void *open(const handle_data_t &handle_data, const sycl::device &dev); @@ -250,7 +259,7 @@ _Effects_: Equivalent to: ---- sycl::device d; sycl::context ctxt = d.get_platform().khr_get_default_context(); -return ipc_memory::get(ptr, ctxt); +return ipc::memory::get(ptr, ctxt); ---- !==== @@ -291,7 +300,7 @@ _Effects_: Equivalent to: ---- sycl::device d; sycl::context ctxt = d.get_platform().khr_get_default_context(); -ipc_memory::put(handle_data, ctxt); +ipc::memory::put(ipc_handle, ctxt); ---- !==== @@ -352,15 +361,14 @@ _Effects_: Equivalent to: [source,c++,indent=2] ---- sycl::context ctxt = dev.get_platform().khr_get_default_context(); -return ipc_memory::put(handle_data, ctxt, dev); +return ipc::memory::open(handle_data, ctxt, dev); ---- !==== a! [source] ---- -void *open(const handle_data_t &handle_data, const sycl::context &ctx, - const sycl::device &dev) +void *open(const handle_data_t &handle_data) ---- !==== @@ -370,15 +378,15 @@ _Effects_: Equivalent to: ---- sycl::device d; sycl::context ctxt = d.get_platform().khr_get_default_context(); -return ipc_memory::open(handle_data, ctxt, d); +return ipc::memory::open(handle_data, ctxt, d); ---- !==== a! [source] ---- -void *open(const handle_data_view_t &handle_data_view, const sycl::context &ctx, - const sycl::device &dev) +void *open(const handle_data_view_t &handle_data_view, + const sycl::context &ctx, const sycl::device &dev) ---- !==== @@ -387,7 +395,7 @@ _Effects_: Equivalent to: [source,c++,indent=2] ---- handle_data_t handle_data{handle_data_view.begin(), handle_data_view.end()}; -return ipc_memory::put(handle_data, ctx, dev); +return ipc::memory::open(handle_data, ctx, dev); ---- !==== @@ -403,15 +411,14 @@ _Effects_: Equivalent to: [source,c++,indent=2] ---- sycl::context ctxt = dev.get_platform().khr_get_default_context(); -return ipc_memory::put(handle_data_view, ctxt, dev); +return ipc::memory::open(handle_data_view, ctxt, dev); ---- !==== a! [source] ---- -void *open(const handle_data_view_t &handle_data_view, const sycl::context &ctx, - const sycl::device &dev) +void *open(const handle_data_view_t &handle_data_view) ---- !==== @@ -421,7 +428,7 @@ _Effects_: Equivalent to: ---- sycl::device d; sycl::context ctxt = d.get_platform().khr_get_default_context(); -return ipc_memory::open(handle_data_view, ctxt, d); +return ipc::memory::open(handle_data_view, ctxt, d); ---- !==== @@ -453,11 +460,661 @@ _Effects_: Equivalent to: ---- sycl::device d; sycl::context ctxt = d.get_platform().khr_get_default_context(); -ipc_memory::close(ptr, ctxt); +ipc::memory::close(ptr, ctxt); +---- + +|==== + +==== New functions and types associated with USM memory (deprecated) + +This extension adds a new type aliases `handle_data_t` and `handle_data_view_t` +as well as a new IPC memory handle type `handle` under the `ipc_memory` +experimental namespace. + +Note, that the `ipc_memory` namespace is deprecated and applications should use +the `ipc::memory` namespace instead. + +``` +namespace sycl::ext::oneapi::experimental { +namespace [[deprecated]] ipc_memory { + +using handle_data_t = std::vector; + +// Requires C++20 +using handle_data_view_t = std::span; + +struct handle { + handle_data_t data() const; + + // Requires C++20 + handle_data_view_t data_view() const; +}; + +} +} +``` + +|==== +a| +[frame=all,grid=none] +!==== +a! +[source] +---- +handle_data_t data() const; +---- +!==== + +_Preconditions:_ The `put` function has not previously been called with this +handle and the USM device memory associated with this handle has not been freed. + +_Returns:_ The handle data associated with the IPC handle object. This data can +be transferred to other processes, but cannot be used to recreate a `handle` +object. + +!==== +a! +[source] +---- +handle_data_view_t data_view() const; ---- +!==== + +_Preconditions:_ The `put` function has not previously been called with this +handle and the USM device memory associated with this handle has not been freed. + +_Returns:_ A `std::span` with a view into the handle data associated with the +IPC handle object. This data can be transferred to other processes, but cannot +be used to recreate a `handle` object. + + +|==== + +Additionally, this extension adds new free functions under the `ipc_memory` +experimental namespace. + +``` +namespace sycl::ext::oneapi::experimental { +namespace [[deprecated]] ipc_memory { + +handle get(void *ptr, const sycl::context &ctx); + +handle get(void *ptr); + +void put(handle &ipc_handle, const sycl::context &ctx); + +void put(handle &ipc_handle); + +void *open(const handle_data_t &handle_data, const sycl::context &ctx, + const sycl::device &dev); + +void *open(const handle_data_t &handle_data, const sycl::device &dev); + +void *open(const handle_data_t &handle_data); + +// Requires C++20 +void *open(const handle_data_view_t &handle_data, const sycl::context &ctx, + const sycl::device &dev); + +// Requires C++20 +void *open(const handle_data_view_t &handle_data, const sycl::device &dev); + +// Requires C++20 +void *open(const handle_data_view_t &handle_data); + +void close(void *ptr, const sycl::context &ctx); + +void close(void *ptr); + +} +} +``` |==== +a| +[frame=all,grid=none] +!==== +a! +[source] +---- +handle get(void *ptr, const sycl::context &ctx) +---- +!==== + +_Preconditions:_ `ptr` is a pointer to USM device memory on some device _D_, and +`ctx` is the same context as `ptr` was allocated against, using the USM device +memory allocation routines. + +_Returns:_ An IPC "handle" to this USM memory allocation. The bytes of this +handle can be transferred to another process on the same system, and the other +process can use the handle to get a pointer to the same USM allocation through a +call to the `open` function. + +_Throws:_ An exception with the `errc::feature_not_supported` error code if +device _D_ does not have `aspect::ext_oneapi_ipc_memory`. + +!==== +a! +[source] +---- +handle get(void *ptr) +---- +!==== + +_Effects_: Equivalent to: + +[source,c++,indent=2] +---- +sycl::device d; +sycl::context ctxt = d.get_platform().khr_get_default_context(); +return ipc_memory::get(ptr, ctxt); +---- + +!==== +a! +[source] +---- +void put(handle &ipc_handle, const sycl::context &ctx) +---- +!==== + +_Preconditions:_ `ipc_handle` is the IPC "handle" to USM device memory that was +returned from a call to `get`. The `put` function has not been previously called on the +handle. + +_Effects:_ Deallocates resources associated with the handle. These resources are +automatically deallocated when the USM device memory is freed, so it is not +strictly necessary to call the `put` function. After the resources associated +with the handle have been deallocated, i.e. through a call to the `put` function +or through freeing the USM device memory, the handle data is invalid and using +it in the `put` and `open` functions will result in undefined behavior. + +[_Note:_ Any pointers retrieved through a call to the `open` function in any +process on the system will still be valid after a call to the `put` function and +must still be freed through calls to the `close` function. +_{endnote}_] + +!==== +a! +[source] +---- +void put(handle &ipc_handle) +---- +!==== + +_Effects_: Equivalent to: + +[source,c++,indent=2] +---- +sycl::device d; +sycl::context ctxt = d.get_platform().khr_get_default_context(); +ipc_memory::put(ipc_handle, ctxt); +---- + +!==== +a! +[source] +---- +void *open(const handle_data_t &handle_data, const sycl::context &ctx, + const sycl::device &dev) +---- +!==== + +_Preconditions:_ `handle_data` is the IPC "handle" to USM device memory that was +returned from a call to the `get` function either in this process or in some +other process on the same system. That USM device memory is accessible on device +`dev`. + +_Returns:_ A pointer to the same USM device memory represented by `handle_data`. +The returned pointer is associated with context `ctx`. It can be used wherever a +USM device pointer for device `dev` and context `ctx` are expected, except it +cannot be passed to `sycl::free`. Instead, use the `close` function to free this +memory pointer. +[_Note:_ The `open` function can be called multiple times on the same handle +within the same process. Each call to the `open` function may return a unique +pointer value even for the same handle, therefore each call to the `open` +function must have a matching call to the `close` function. +_{endnote}_] + +[_Note:_ The pointer returned from a call to the `open` function is no longer +valid if the associated USM device memory is freed through a call to the +`sycl::free` function. +_{endnote}_] + +[_Note:_ For sub-page allocations, the implementation may need to share the +entire page with the other process. As a result, the other process may have +access to adjacent USM memory that happens to share the same page. Applications +should be aware of the potential security implications if the USM memory +contains sensitive information, and the other application isn't trusted. +_{endnote}_] + +_Throws:_ + + * An exception with the `errc::feature_not_supported` error code if device + `dev` does not have `aspect::ext_oneapi_ipc_memory`. + * An exception with the `errc::invalid` error code if the handle data + `handle_data` has an unexpected number of bytes. + +!==== +a! +[source] +---- +void *open(const handle_data_t &handle_data, const sycl::device &dev) +---- +!==== + +_Effects_: Equivalent to: + +[source,c++,indent=2] +---- +sycl::context ctxt = dev.get_platform().khr_get_default_context(); +return ipc_memory::open(handle_data, ctxt, dev); +---- + +!==== +a! +[source] +---- +void *open(const handle_data_t &handle_data) +---- +!==== + +_Effects_: Equivalent to: + +[source,c++,indent=2] +---- +sycl::device d; +sycl::context ctxt = d.get_platform().khr_get_default_context(); +return ipc_memory::open(handle_data, ctxt, d); +---- + +!==== +a! +[source] +---- +void *open(const handle_data_view_t &handle_data_view, const sycl::context &ctx, + const sycl::device &dev) +---- +!==== + +_Effects_: Equivalent to: + +[source,c++,indent=2] +---- +handle_data_t handle_data{handle_data_view.begin(), handle_data_view.end()}; +return ipc_memory::open(handle_data, ctx, dev); +---- + +!==== +a! +[source] +---- +void *open(const handle_data_view_t &handle_data_view, const sycl::device &dev) +---- +!==== + +_Effects_: Equivalent to: + +[source,c++,indent=2] +---- +sycl::context ctxt = dev.get_platform().khr_get_default_context(); +return ipc_memory::open(handle_data_view, ctxt, dev); +---- + +!==== +a! +[source] +---- +void *open(const handle_data_view_t &handle_data_view) +---- +!==== + +_Effects_: Equivalent to: + +[source,c++,indent=2] +---- +sycl::device d; +sycl::context ctxt = d.get_platform().khr_get_default_context(); +return ipc_memory::open(handle_data_view, ctxt, d); +---- + +!==== +a! +[source] +---- +void close(void *ptr, const sycl::context &ctx) +---- +!==== + +_Precondition:_ `ptr` was previously returned from a call to the `open` function +in this same process, where `ctx` was passed as the context. This `ptr` value +has not yet been closed by calling the `close` function. + +_Effects:_ Closes a device USM pointer previously returned by a call to +the `open` function. + +!==== +a! +[source] +---- +void close(void *ptr) +---- +!==== + +_Effects_: Equivalent to: + +[source,c++,indent=2] +---- +sycl::device d; +sycl::context ctxt = d.get_platform().khr_get_default_context(); +ipc_memory::close(ptr, ctxt); +---- + +|==== + +=== Inter-process exchange of physical memory + +This extension allows for the exchange of a physical memory handle between +processes. A given physical memory handle can be mapped into a virtual address +space in multiple processes, so the processes can access the same device +memory. + +The following aspect and set of free functions are specific to +physical memory handle sharing feature. + +The physical memory is represented by the +`sycl::ext::oneapi::experimental::physical_mem` object instance +(from the link:sycl_ext_oneapi_virtual_mem.asciidoc[ +sycl_ext_oneapi_virtual_mem] extension). + +==== Extension to `enum class aspect` for physical memory + +[source] +---- +namespace sycl { +enum class aspect { + ... + ext_oneapi_ipc_physical_memory +} +} +---- + +If a SYCL device has this aspect, that device supports the functions in the +`ipc::physical_memory` namespace specified in the following section. + +==== New property for the physical_mem constructor + +This extension adds the following property, which can be passed to the +physical_mem object constructor. + +[source,c++] +---- +namespace sycl::ext::oneapi::experimental { + +struct enable_ipc { + enable_ipc(bool enable = true); +}; +using enable_ipc_key = enable_ipc; + +} // namespace sycl::ext::oneapi::experimental +---- + +This property controls whether the physical memory object can be shared across +processes. + +_Effects:_ Creates a new `enable_ipc` property with a boolean value +indicating whether the physical memory object can be shared across processes. + +_Remarks:_ `physical_mem` constructor throws an exception with the +`errc::feature_not_supported` error code if inter-process sharing is enabled +via `enable_ipc` and the device associated with `physical_mem` does not have +the aspect `aspect::ext_oneapi_ipc_physical_memory`. + +==== New physical_mem member function + +This extension adds a new physical_mem type member function, which can be used +to check if a given physical memory allocation can be passed between processes. + +[source,c++] +---- +namespace sycl::ext::oneapi::experimental { + +class physical_mem { + // ... + bool ipc_enabled() const; +}; + +} // namespace sycl::ext::oneapi::experimental +---- + +_Returns:_ True if the physical_mem object was created with inter-process +sharing enabled. + +==== New functions associated with physical memory + +This extension adds new free functions under the `ipc::physical_memory` experimental +namespace. + +``` +namespace sycl::ext::oneapi::experimental::ipc::physical_memory { + +handle get(const physical_mem &phys_mem); + +void put(handle &ipc_handle, const sycl::context &ctx); + +void put(handle &ipc_handle); + +physical_mem open(const handle_data_t &handle_data, const sycl::context &ctx, + const sycl::device &dev); + +physical_mem open(const handle_data_t &handle_data, const sycl::device &dev); + +physical_mem open(const handle_data_t &handle_data); + +// Requires C++20 +physical_mem open(const handle_data_view_t &handle_data_view, + const sycl::context &ctx, const sycl::device &dev); + +// Requires C++20 +physical_mem open(const handle_data_view_t &handle_data_view, + const sycl::device &dev); + +// Requires C++20 +physical_mem open(const handle_data_view_t &handle_data_view); + +} +``` + +|==== +a| +[frame=all,grid=none] +!==== +a! +[source] +---- +handle get(const physical_mem &phys_mem) +---- +!==== + +_Returns:_ An IPC "handle" to this physical memory object. The bytes of this +handle can be transferred to another process on the same system, and the other +process can use the handle to get a physical memory object representing the +same memory allocation through a call to the `open` function. + +_Throws:_ An exception with the `errc::invalid` error code if `phys_mem` was +not created with inter-process sharing enabled via the `enable_ipc` property. + +!==== +a! +[source] +---- +void put(handle &ipc_handle, const sycl::context &ctx) +---- +!==== + +_Preconditions:_ + + * `ipc_handle` is the IPC "handle" to physical memory that was returned from + a call to `get`. The `put` function has not been previously called on the + handle. + * `ctx` is a context associated with the `physical_mem` object, which was + passed to the `get` function call that produced `ipc_handle`. + +_Effects:_ Deallocates resources associated with the handle. These resources are +automatically deallocated when the physical memory object is destroyed, so it +is not strictly necessary to call the `put` function. After the resources +associated with the handle have been deallocated, i.e. through a call to the +`put` function or through destruction of the physical memory object, the handle +data is invalid and using it in the `put` and `open` functions will result in +undefined behavior. + +[_Note:_ Any objects retrieved through a call to the `open` function in any +process on the system will still be valid after a call to the `put` function. +_{endnote}_] + +!==== +a! +[source] +---- +void put(handle &ipc_handle) +---- +!==== + +_Effects_: Equivalent to: + +[source,c++,indent=2] +---- +sycl::device d; +sycl::context ctxt = d.get_platform().khr_get_default_context(); +ipc::physical_memory::put(ipc_handle, ctxt); +---- + +!==== +a! +[source] +---- +physical_mem open(const handle_data_t &handle_data, const sycl::context &ctx, + const sycl::device &dev) +---- +!==== + +_Preconditions:_ + + * `handle_data` is the IPC "handle" to a physical memory object + that was returned from a call to the `get` function either in this process + or in some other process on the same system. + * `dev` is the same device that was associated with the `physical_mem` object + passed to the `get` function call that produced `handle_data`. + +_Returns:_ A physical memory object represented by `handle_data`. The returned +object is associated with context `ctx` and device `dev`. + +[_Note:_ The `open` function can be called multiple times on the same handle +within the same process. Each call to the `open` function may return a unique +physical memory object even for the same handle. +_{endnote}_] + +[_Note:_ The physical memory allocation associated with the object returned +from a call to the `open` function is valid through the entire lifetime of the +physical memory object. The physical memory is released once all the physical +memory objects associated with that allocation are destroyed. +_{endnote}_] + +_Throws:_ + + * An exception with the `errc::feature_not_supported` error code if device + `dev` does not have `aspect::ext_oneapi_ipc_physical_memory`. + * An exception with the `errc::invalid` error code if the handle data + `handle_data` has an unexpected number of bytes. + * An exception with the `errc::invalid` error code if `ctx` does not contain + `dev`. + +!==== +a! +[source] +---- +physical_mem open(const handle_data_t &handle_data, const sycl::device &dev) +---- +!==== + +_Effects:_ Equivalent to: + +[source,c++,indent=2] +---- +sycl::context ctxt = dev.get_platform().khr_get_default_context(); +return ipc::physical_memory::open(handle_data, ctxt, dev); +---- + +!==== +a! +[source] +---- +physical_mem open(const handle_data_t &handle_data) +---- +!==== + +_Effects:_ Equivalent to: + +[source,c++,indent=2] +---- +sycl::device d; +sycl::context ctxt = d.get_platform().khr_get_default_context(); +return ipc::physical_memory::open(handle_data, ctxt, d); +---- + +!==== +a! +[source] +---- +physical_mem open(const handle_data_view_t &handle_data_view, + const sycl::context &ctx, const sycl::device &dev) +---- +!==== + +_Effects:_ Equivalent to: + +[source,c++,indent=2] +---- +handle_data_t handle_data{handle_data_view.begin(), handle_data_view.end()}; +return ipc::physical_memory::open(handle_data, ctx, dev); +---- + +!==== +a! +[source] +---- +physical_mem open(const handle_data_view_t &handle_data_view, + const sycl::device &dev) +---- +!==== + +_Effects:_ Equivalent to: + +[source,c++,indent=2] +---- +sycl::context ctxt = dev.get_platform().khr_get_default_context(); +return ipc::physical_memory::open(handle_data_view, ctxt, dev); +---- + +!==== +a! +[source] +---- +physical_mem open(const handle_data_view_t &handle_data_view) +---- +!==== + +_Effects:_ Equivalent to: + +[source,c++,indent=2] +---- +sycl::device d; +sycl::context ctxt = d.get_platform().khr_get_default_context(); +return ipc::physical_memory::open(handle_data_view, ctxt, d); +---- + +|==== == Issues diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_virtual_mem.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_virtual_mem.asciidoc index 0432f01f7168..e86df0a6c221 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_virtual_mem.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_virtual_mem.asciidoc @@ -41,6 +41,9 @@ This extension is written against the SYCL 2020 revision 8 specification. All references below to the "core SYCL specification" or to section numbers in the SYCL specification refer to that revision. +This extension also depends on the following other SYCL extensions: + +* link:sycl_ext_oneapi_properties.asciidoc[sycl_ext_oneapi_properties] == Status @@ -242,8 +245,11 @@ enum class address_access_mode : /*unspecified*/ { class physical_mem { public: - physical_mem(const device &syclDevice, const context &syclContext, size_t numBytes); - physical_mem(const queue &syclQueue, size_t numBytes); + template + physical_mem(const device &syclDevice, const context &syclContext, size_t numBytes, + PropertyListT props = {}); + template + physical_mem(const queue &syclQueue, size_t numBytes, PropertyListT props = {}); /* -- common interface members -- */ @@ -265,7 +271,9 @@ public: |============================ |Member function |Description -|`physical_mem(const device &syclDevice, const context &syclContext, size_t numBytes)` | +|`template + physical_mem(const device &syclDevice, const context &syclContext, size_t numBytes, + PropertyListT props = {})` | Constructs a `physical_mem` instance using the `syclDevice` provided. This device must either be contained by `syclContext` or it must be a descendent device of some device that is contained by that context, otherwise this function @@ -281,8 +289,12 @@ an exception with `errc::feature_not_supported`. If the constructor is unable to allocate the required memory on `syclDevice`, the call throws an exception with `errc::memory_allocation`. -|`physical_mem(const queue &syclQueue, size_t numBytes)` | -Same as `physical_mem(syclQueue.get_device(), syclQueue.get_context(), numBytes)`. +`props` can optionally be used to provide the physical memory properties. + +|`template + physical_mem(const queue &syclQueue, size_t numBytes, PropertyListT props = {})` | +Same as `physical_mem(syclQueue.get_device(), syclQueue.get_context(), numBytes, +props)`. |`void *map(uintptr_t ptr, size_t numBytes, address_access_mode mode, size_t offset = 0)` | Maps a virtual memory range, specified by `ptr` and `numBytes`, to the physical From 0b4c68aca3fea88765f31d7604dd8fa01da625a9 Mon Sep 17 00:00:00 2001 From: Lukasz Dorau Date: Wed, 20 May 2026 14:41:22 +0000 Subject: [PATCH 02/17] [UR] Add IPC spec for physical_mem objects Add four new function specs to exp-inter-process-communication.yml and update registry.yml to support IPC operations on physical memory objects: - urIPCGetPhysMemHandleExp: export an IPC handle for a physical_mem - urIPCPutPhysMemHandleExp: release the exported IPC handle - urIPCOpenPhysMemHandleExp: import an IPC handle into a physical_mem - urIPCClosePhysMemHandleExp: close the imported physical_mem Signed-off-by: Lukasz Dorau --- .../core/exp-inter-process-communication.yml | 104 ++++++++++++++++++ unified-runtime/scripts/core/registry.yml | 14 ++- 2 files changed, 117 insertions(+), 1 deletion(-) diff --git a/unified-runtime/scripts/core/exp-inter-process-communication.yml b/unified-runtime/scripts/core/exp-inter-process-communication.yml index e5019bb40836..312dba7919ad 100644 --- a/unified-runtime/scripts/core/exp-inter-process-communication.yml +++ b/unified-runtime/scripts/core/exp-inter-process-communication.yml @@ -126,3 +126,107 @@ returns: - "`NULL == pMem`" - $X_RESULT_ERROR_OUT_OF_HOST_MEMORY - $X_RESULT_ERROR_OUT_OF_RESOURCES +--- #-------------------------------------------------------------------------- +type: function +desc: "Gets an inter-process handle for a physical memory object" +class: $xIPC +name: GetPhysMemHandleExp +ordinal: "0" +params: + - type: $x_context_handle_t + name: hContext + desc: "[in] handle of the context object" + - type: $x_physical_mem_handle_t + name: hPhysMem + desc: "[in] handle of the physical memory object" + - type: void** + name: ppIPCPhysMemHandleData + desc: "[out][optional] a pointer to the IPC physical memory handle data" + - type: size_t* + name: pIPCPhysMemHandleDataSizeRet + desc: "[out][optional] size of the resulting IPC physical memory handle data" +returns: + - $X_RESULT_ERROR_INVALID_CONTEXT + - $X_RESULT_ERROR_INVALID_NULL_HANDLE: + - "`NULL == hContext`" + - "`NULL == hPhysMem`" + - $X_RESULT_ERROR_INVALID_NULL_POINTER: + - "`NULL == ppIPCPhysMemHandleData`" + - "`NULL == pIPCPhysMemHandleDataSizeRet`" + - $X_RESULT_ERROR_OUT_OF_HOST_MEMORY + - $X_RESULT_ERROR_OUT_OF_RESOURCES +--- #-------------------------------------------------------------------------- +type: function +desc: "Releases an inter-process physical memory handle" +class: $xIPC +name: PutPhysMemHandleExp +ordinal: "0" +params: + - type: $x_context_handle_t + name: hContext + desc: "[in] handle of the context object" + - type: void* + name: pIPCPhysMemHandleData + desc: "[in] a pointer to the IPC physical memory handle data" +returns: + - $X_RESULT_ERROR_INVALID_CONTEXT + - $X_RESULT_ERROR_INVALID_NULL_HANDLE: + - "`NULL == hContext`" + - $X_RESULT_ERROR_INVALID_NULL_POINTER: + - "`NULL == pIPCPhysMemHandleData`" + - $X_RESULT_ERROR_OUT_OF_HOST_MEMORY + - $X_RESULT_ERROR_OUT_OF_RESOURCES +--- #-------------------------------------------------------------------------- +type: function +desc: "Opens an inter-process physical memory handle to get the corresponding physical memory object" +class: $xIPC +name: OpenPhysMemHandleExp +ordinal: "0" +params: + - type: $x_context_handle_t + name: hContext + desc: "[in] handle of the context object" + - type: $x_device_handle_t + name: hDevice + desc: "[in] handle of the device object the physical memory was allocated on" + - type: void * + name: pIPCPhysMemHandleData + desc: "[in] the IPC physical memory handle data" + - type: size_t + name: ipcPhysMemHandleDataSize + desc: "[in] size of the IPC physical memory handle data" + - type: $x_physical_mem_handle_t* + name: phPhysMem + desc: "[out] pointer to the physical memory handle" +returns: + - $X_RESULT_ERROR_INVALID_CONTEXT + - $X_RESULT_ERROR_INVALID_NULL_HANDLE: + - "`NULL == hContext`" + - "`NULL == hDevice`" + - $X_RESULT_ERROR_INVALID_NULL_POINTER: + - "`NULL == pIPCPhysMemHandleData`" + - "`NULL == phPhysMem`" + - $X_RESULT_ERROR_INVALID_VALUE: + - "ipcPhysMemHandleDataSize is not the same as the size of IPC physical memory handle data" + - $X_RESULT_ERROR_OUT_OF_HOST_MEMORY + - $X_RESULT_ERROR_OUT_OF_RESOURCES +--- #-------------------------------------------------------------------------- +type: function +desc: "Closes an inter-process physical memory handle" +class: $xIPC +name: ClosePhysMemHandleExp +ordinal: "0" +params: + - type: $x_context_handle_t + name: hContext + desc: "[in] handle of the context object" + - type: $x_physical_mem_handle_t + name: hPhysMem + desc: "[in] physical memory handle opened through urIPCOpenPhysMemHandleExp" +returns: + - $X_RESULT_ERROR_INVALID_CONTEXT + - $X_RESULT_ERROR_INVALID_NULL_HANDLE: + - "`NULL == hContext`" + - "`NULL == hPhysMem`" + - $X_RESULT_ERROR_OUT_OF_HOST_MEMORY + - $X_RESULT_ERROR_OUT_OF_RESOURCES diff --git a/unified-runtime/scripts/core/registry.yml b/unified-runtime/scripts/core/registry.yml index 6c039b80e41d..b83ae82f2f12 100644 --- a/unified-runtime/scripts/core/registry.yml +++ b/unified-runtime/scripts/core/registry.yml @@ -724,7 +724,19 @@ etors: - name: QUEUE_GET_GRAPH_EXP desc: Enumerator for $xQueueGetGraphExp value: '314' -max_id: '314' +- name: IPC_GET_PHYS_MEM_HANDLE_EXP + desc: Enumerator for $xIPCGetPhysMemHandleExp + value: '315' +- name: IPC_PUT_PHYS_MEM_HANDLE_EXP + desc: Enumerator for $xIPCPutPhysMemHandleExp + value: '316' +- name: IPC_OPEN_PHYS_MEM_HANDLE_EXP + desc: Enumerator for $xIPCOpenPhysMemHandleExp + value: '317' +- name: IPC_CLOSE_PHYS_MEM_HANDLE_EXP + desc: Enumerator for $xIPCClosePhysMemHandleExp + value: '318' +max_id: '318' --- type: enum desc: Defines structure types From 1b7b18d7ba738ef36d2fb8333a419065dac47692 Mon Sep 17 00:00:00 2001 From: Lukasz Dorau Date: Wed, 20 May 2026 14:52:20 +0000 Subject: [PATCH 03/17] [UR] Add generated files for IPC physical_mem API Add auto-generated files based on the IPC physical_mem spec changes (urIPCGetPhysMemHandleExp, urIPCPutPhysMemHandleExp, urIPCOpenPhysMemHandleExp, urIPCClosePhysMemHandleExp): - ur_api.h, ur_api_funcs.def, ur_ddi.h: function declarations - ur_print.h/hpp, ur_print.cpp: printing support - ur_ldrddi.cpp, ur_libapi.cpp, ur_api.cpp: loader implementations - ur_trcddi.cpp, ur_valddi.cpp: tracing and validation layers - ur_mockddi.cpp: mock adapter - loader.def.in, loader.map.in: loader symbol exports - level_zero ur_interface_loader.cpp/.hpp: Level-Zero DDI table entries - opencl common.hpp: updated generated include Signed-off-by: Lukasz Dorau --- .../include/unified-runtime/ur_api.h | 151 +++++++++++++ .../include/unified-runtime/ur_api_funcs.def | 4 + .../include/unified-runtime/ur_ddi.h | 25 +++ .../include/unified-runtime/ur_print.h | 40 ++++ .../include/unified-runtime/ur_print.hpp | 133 ++++++++++++ .../level_zero/ur_interface_loader.cpp | 6 + .../level_zero/ur_interface_loader.hpp | 13 ++ .../source/adapters/mock/ur_mockddi.cpp | 200 ++++++++++++++++++ .../source/adapters/opencl/common.hpp | 24 +-- .../loader/layers/tracing/ur_trcddi.cpp | 186 ++++++++++++++++ .../loader/layers/validation/ur_valddi.cpp | 184 ++++++++++++++++ unified-runtime/source/loader/loader.def.in | 8 + unified-runtime/source/loader/loader.map.in | 8 + unified-runtime/source/loader/ur_ldrddi.cpp | 89 ++++++++ unified-runtime/source/loader/ur_libapi.cpp | 140 ++++++++++++ unified-runtime/source/loader/ur_print.cpp | 32 +++ unified-runtime/source/ur_api.cpp | 114 ++++++++++ 17 files changed, 1345 insertions(+), 12 deletions(-) diff --git a/unified-runtime/include/unified-runtime/ur_api.h b/unified-runtime/include/unified-runtime/ur_api.h index ce889941d39c..826feace5c79 100644 --- a/unified-runtime/include/unified-runtime/ur_api.h +++ b/unified-runtime/include/unified-runtime/ur_api.h @@ -510,6 +510,14 @@ typedef enum ur_function_t { UR_FUNCTION_USM_HOST_ALLOC_UNREGISTER_EXP = 313, /// Enumerator for ::urQueueGetGraphExp UR_FUNCTION_QUEUE_GET_GRAPH_EXP = 314, + /// Enumerator for ::urIPCGetPhysMemHandleExp + UR_FUNCTION_IPC_GET_PHYS_MEM_HANDLE_EXP = 315, + /// Enumerator for ::urIPCPutPhysMemHandleExp + UR_FUNCTION_IPC_PUT_PHYS_MEM_HANDLE_EXP = 316, + /// Enumerator for ::urIPCOpenPhysMemHandleExp + UR_FUNCTION_IPC_OPEN_PHYS_MEM_HANDLE_EXP = 317, + /// Enumerator for ::urIPCClosePhysMemHandleExp + UR_FUNCTION_IPC_CLOSE_PHYS_MEM_HANDLE_EXP = 318, /// @cond UR_FUNCTION_FORCE_UINT32 = 0x7fffffff /// @endcond @@ -11240,6 +11248,108 @@ UR_APIEXPORT ur_result_t UR_APICALL urIPCCloseMemHandleExp( /// [in] pointer to device USM memory opened through urIPCOpenMemHandleExp void *pMem); +/////////////////////////////////////////////////////////////////////////////// +/// @brief Gets an inter-process handle for a physical memory object +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hContext` +/// + `NULL == hPhysMem` +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == ppIPCPhysMemHandleData` +/// + `NULL == pIPCPhysMemHandleDataSizeRet` +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +UR_APIEXPORT ur_result_t UR_APICALL urIPCGetPhysMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the physical memory object + ur_physical_mem_handle_t hPhysMem, + /// [out][optional] a pointer to the IPC physical memory handle data + void **ppIPCPhysMemHandleData, + /// [out][optional] size of the resulting IPC physical memory handle data + size_t *pIPCPhysMemHandleDataSizeRet); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Releases an inter-process physical memory handle +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hContext` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pIPCPhysMemHandleData` +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +UR_APIEXPORT ur_result_t UR_APICALL urIPCPutPhysMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] a pointer to the IPC physical memory handle data + void *pIPCPhysMemHandleData); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Opens an inter-process physical memory handle to get the +/// corresponding +/// physical memory object +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hContext` +/// + `NULL == hDevice` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == phPhysMem` +/// + `NULL == pIPCPhysMemHandleData` +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +/// - ::UR_RESULT_ERROR_INVALID_VALUE +/// + ipcPhysMemHandleDataSize is not the same as the size of IPC +/// physical memory handle data +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +UR_APIEXPORT ur_result_t UR_APICALL urIPCOpenPhysMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object the physical memory was allocated on + ur_device_handle_t hDevice, + /// [in] the IPC physical memory handle data + void *pIPCPhysMemHandleData, + /// [in] size of the IPC physical memory handle data + size_t ipcPhysMemHandleDataSize, + /// [out] pointer to the physical memory handle + ur_physical_mem_handle_t *phPhysMem); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Closes an inter-process physical memory handle +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hContext` +/// + `NULL == hPhysMem` +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +UR_APIEXPORT ur_result_t UR_APICALL urIPCClosePhysMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] physical memory handle opened through urIPCOpenPhysMemHandleExp + ur_physical_mem_handle_t hPhysMem); + #if !defined(__GNUC__) #pragma endregion #endif @@ -16464,6 +16574,47 @@ typedef struct ur_ipc_close_mem_handle_exp_params_t { void **ppMem; } ur_ipc_close_mem_handle_exp_params_t; +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function parameters for urIPCGetPhysMemHandleExp +/// @details Each entry is a pointer to the parameter passed to the function; +/// allowing the callback the ability to modify the parameter's value +typedef struct ur_ipc_get_phys_mem_handle_exp_params_t { + ur_context_handle_t *phContext; + ur_physical_mem_handle_t *phPhysMem; + void ***pppIPCPhysMemHandleData; + size_t **ppIPCPhysMemHandleDataSizeRet; +} ur_ipc_get_phys_mem_handle_exp_params_t; + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function parameters for urIPCPutPhysMemHandleExp +/// @details Each entry is a pointer to the parameter passed to the function; +/// allowing the callback the ability to modify the parameter's value +typedef struct ur_ipc_put_phys_mem_handle_exp_params_t { + ur_context_handle_t *phContext; + void **ppIPCPhysMemHandleData; +} ur_ipc_put_phys_mem_handle_exp_params_t; + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function parameters for urIPCOpenPhysMemHandleExp +/// @details Each entry is a pointer to the parameter passed to the function; +/// allowing the callback the ability to modify the parameter's value +typedef struct ur_ipc_open_phys_mem_handle_exp_params_t { + ur_context_handle_t *phContext; + ur_device_handle_t *phDevice; + void **ppIPCPhysMemHandleData; + size_t *pipcPhysMemHandleDataSize; + ur_physical_mem_handle_t **pphPhysMem; +} ur_ipc_open_phys_mem_handle_exp_params_t; + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function parameters for urIPCClosePhysMemHandleExp +/// @details Each entry is a pointer to the parameter passed to the function; +/// allowing the callback the ability to modify the parameter's value +typedef struct ur_ipc_close_phys_mem_handle_exp_params_t { + ur_context_handle_t *phContext; + ur_physical_mem_handle_t *phPhysMem; +} ur_ipc_close_phys_mem_handle_exp_params_t; + /////////////////////////////////////////////////////////////////////////////// /// @brief Function parameters for urMemoryExportAllocExportableMemoryExp /// @details Each entry is a pointer to the parameter passed to the function; diff --git a/unified-runtime/include/unified-runtime/ur_api_funcs.def b/unified-runtime/include/unified-runtime/ur_api_funcs.def index c73c88213600..84afd9752628 100644 --- a/unified-runtime/include/unified-runtime/ur_api_funcs.def +++ b/unified-runtime/include/unified-runtime/ur_api_funcs.def @@ -222,6 +222,10 @@ _UR_API(urIPCGetMemHandleExp) _UR_API(urIPCPutMemHandleExp) _UR_API(urIPCOpenMemHandleExp) _UR_API(urIPCCloseMemHandleExp) +_UR_API(urIPCGetPhysMemHandleExp) +_UR_API(urIPCPutPhysMemHandleExp) +_UR_API(urIPCOpenPhysMemHandleExp) +_UR_API(urIPCClosePhysMemHandleExp) _UR_API(urMemoryExportAllocExportableMemoryExp) _UR_API(urMemoryExportFreeExportableMemoryExp) _UR_API(urMemoryExportExportMemoryHandleExp) diff --git a/unified-runtime/include/unified-runtime/ur_ddi.h b/unified-runtime/include/unified-runtime/ur_ddi.h index 2fda5130c0ed..1b0c64264714 100644 --- a/unified-runtime/include/unified-runtime/ur_ddi.h +++ b/unified-runtime/include/unified-runtime/ur_ddi.h @@ -1944,6 +1944,27 @@ typedef ur_result_t(UR_APICALL *ur_pfnIPCOpenMemHandleExp_t)( typedef ur_result_t(UR_APICALL *ur_pfnIPCCloseMemHandleExp_t)( ur_context_handle_t, void *); +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function-pointer for urIPCGetPhysMemHandleExp +typedef ur_result_t(UR_APICALL *ur_pfnIPCGetPhysMemHandleExp_t)( + ur_context_handle_t, ur_physical_mem_handle_t, void **, size_t *); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function-pointer for urIPCPutPhysMemHandleExp +typedef ur_result_t(UR_APICALL *ur_pfnIPCPutPhysMemHandleExp_t)( + ur_context_handle_t, void *); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function-pointer for urIPCOpenPhysMemHandleExp +typedef ur_result_t(UR_APICALL *ur_pfnIPCOpenPhysMemHandleExp_t)( + ur_context_handle_t, ur_device_handle_t, void *, size_t, + ur_physical_mem_handle_t *); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function-pointer for urIPCClosePhysMemHandleExp +typedef ur_result_t(UR_APICALL *ur_pfnIPCClosePhysMemHandleExp_t)( + ur_context_handle_t, ur_physical_mem_handle_t); + /////////////////////////////////////////////////////////////////////////////// /// @brief Table of IPCExp functions pointers typedef struct ur_ipc_exp_dditable_t { @@ -1951,6 +1972,10 @@ typedef struct ur_ipc_exp_dditable_t { ur_pfnIPCPutMemHandleExp_t pfnPutMemHandleExp; ur_pfnIPCOpenMemHandleExp_t pfnOpenMemHandleExp; ur_pfnIPCCloseMemHandleExp_t pfnCloseMemHandleExp; + ur_pfnIPCGetPhysMemHandleExp_t pfnGetPhysMemHandleExp; + ur_pfnIPCPutPhysMemHandleExp_t pfnPutPhysMemHandleExp; + ur_pfnIPCOpenPhysMemHandleExp_t pfnOpenPhysMemHandleExp; + ur_pfnIPCClosePhysMemHandleExp_t pfnClosePhysMemHandleExp; } ur_ipc_exp_dditable_t; /////////////////////////////////////////////////////////////////////////////// diff --git a/unified-runtime/include/unified-runtime/ur_print.h b/unified-runtime/include/unified-runtime/ur_print.h index f309275767a0..091239fcbfef 100644 --- a/unified-runtime/include/unified-runtime/ur_print.h +++ b/unified-runtime/include/unified-runtime/ur_print.h @@ -3775,6 +3775,46 @@ UR_APIEXPORT ur_result_t UR_APICALL urPrintIpcCloseMemHandleExpParams( const struct ur_ipc_close_mem_handle_exp_params_t *params, char *buffer, const size_t buff_size, size_t *out_size); +/////////////////////////////////////////////////////////////////////////////// +/// @brief Print ur_ipc_get_phys_mem_handle_exp_params_t struct +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_INVALID_SIZE +/// - `buff_size < out_size` +UR_APIEXPORT ur_result_t UR_APICALL urPrintIpcGetPhysMemHandleExpParams( + const struct ur_ipc_get_phys_mem_handle_exp_params_t *params, char *buffer, + const size_t buff_size, size_t *out_size); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Print ur_ipc_put_phys_mem_handle_exp_params_t struct +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_INVALID_SIZE +/// - `buff_size < out_size` +UR_APIEXPORT ur_result_t UR_APICALL urPrintIpcPutPhysMemHandleExpParams( + const struct ur_ipc_put_phys_mem_handle_exp_params_t *params, char *buffer, + const size_t buff_size, size_t *out_size); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Print ur_ipc_open_phys_mem_handle_exp_params_t struct +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_INVALID_SIZE +/// - `buff_size < out_size` +UR_APIEXPORT ur_result_t UR_APICALL urPrintIpcOpenPhysMemHandleExpParams( + const struct ur_ipc_open_phys_mem_handle_exp_params_t *params, char *buffer, + const size_t buff_size, size_t *out_size); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Print ur_ipc_close_phys_mem_handle_exp_params_t struct +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_INVALID_SIZE +/// - `buff_size < out_size` +UR_APIEXPORT ur_result_t UR_APICALL urPrintIpcClosePhysMemHandleExpParams( + const struct ur_ipc_close_phys_mem_handle_exp_params_t *params, + char *buffer, const size_t buff_size, size_t *out_size); + /////////////////////////////////////////////////////////////////////////////// /// @brief Print ur_memory_export_alloc_exportable_memory_exp_params_t struct /// @returns diff --git a/unified-runtime/include/unified-runtime/ur_print.hpp b/unified-runtime/include/unified-runtime/ur_print.hpp index fba0340dea51..d0728babf7a6 100644 --- a/unified-runtime/include/unified-runtime/ur_print.hpp +++ b/unified-runtime/include/unified-runtime/ur_print.hpp @@ -1377,6 +1377,18 @@ inline std::ostream &operator<<(std::ostream &os, enum ur_function_t value) { case UR_FUNCTION_QUEUE_GET_GRAPH_EXP: os << "UR_FUNCTION_QUEUE_GET_GRAPH_EXP"; break; + case UR_FUNCTION_IPC_GET_PHYS_MEM_HANDLE_EXP: + os << "UR_FUNCTION_IPC_GET_PHYS_MEM_HANDLE_EXP"; + break; + case UR_FUNCTION_IPC_PUT_PHYS_MEM_HANDLE_EXP: + os << "UR_FUNCTION_IPC_PUT_PHYS_MEM_HANDLE_EXP"; + break; + case UR_FUNCTION_IPC_OPEN_PHYS_MEM_HANDLE_EXP: + os << "UR_FUNCTION_IPC_OPEN_PHYS_MEM_HANDLE_EXP"; + break; + case UR_FUNCTION_IPC_CLOSE_PHYS_MEM_HANDLE_EXP: + os << "UR_FUNCTION_IPC_CLOSE_PHYS_MEM_HANDLE_EXP"; + break; default: os << "unknown enumerator"; break; @@ -21694,6 +21706,115 @@ operator<<(std::ostream &os, return os; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Print operator for the ur_ipc_get_phys_mem_handle_exp_params_t type +/// @returns +/// std::ostream & +inline std::ostream & +operator<<(std::ostream &os, + [[maybe_unused]] const struct ur_ipc_get_phys_mem_handle_exp_params_t + *params) { + + os << ".hContext = "; + + ur::details::printPtr(os, *(params->phContext)); + + os << ", "; + os << ".hPhysMem = "; + + ur::details::printPtr(os, *(params->phPhysMem)); + + os << ", "; + os << ".ppIPCPhysMemHandleData = "; + + ur::details::printPtr(os, *(params->pppIPCPhysMemHandleData)); + + os << ", "; + os << ".pIPCPhysMemHandleDataSizeRet = "; + + ur::details::printPtr(os, *(params->ppIPCPhysMemHandleDataSizeRet)); + + return os; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Print operator for the ur_ipc_put_phys_mem_handle_exp_params_t type +/// @returns +/// std::ostream & +inline std::ostream & +operator<<(std::ostream &os, + [[maybe_unused]] const struct ur_ipc_put_phys_mem_handle_exp_params_t + *params) { + + os << ".hContext = "; + + ur::details::printPtr(os, *(params->phContext)); + + os << ", "; + os << ".pIPCPhysMemHandleData = "; + + ur::details::printPtr(os, *(params->ppIPCPhysMemHandleData)); + + return os; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Print operator for the ur_ipc_open_phys_mem_handle_exp_params_t type +/// @returns +/// std::ostream & +inline std::ostream &operator<<( + std::ostream &os, + [[maybe_unused]] const struct ur_ipc_open_phys_mem_handle_exp_params_t + *params) { + + os << ".hContext = "; + + ur::details::printPtr(os, *(params->phContext)); + + os << ", "; + os << ".hDevice = "; + + ur::details::printPtr(os, *(params->phDevice)); + + os << ", "; + os << ".pIPCPhysMemHandleData = "; + + os << *(params->ppIPCPhysMemHandleData); + + os << ", "; + os << ".ipcPhysMemHandleDataSize = "; + + os << *(params->pipcPhysMemHandleDataSize); + + os << ", "; + os << ".phPhysMem = "; + + ur::details::printPtr(os, *(params->pphPhysMem)); + + return os; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Print operator for the ur_ipc_close_phys_mem_handle_exp_params_t type +/// @returns +/// std::ostream & +inline std::ostream &operator<<( + std::ostream &os, + [[maybe_unused]] const struct ur_ipc_close_phys_mem_handle_exp_params_t + *params) { + + os << ".hContext = "; + + ur::details::printPtr(os, *(params->phContext)); + + os << ", "; + os << ".hPhysMem = "; + + ur::details::printPtr(os, *(params->phPhysMem)); + + return os; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Print operator for the /// ur_memory_export_alloc_exportable_memory_exp_params_t type @@ -23247,6 +23368,18 @@ inline ur_result_t UR_APICALL printFunctionParams(std::ostream &os, case UR_FUNCTION_IPC_CLOSE_MEM_HANDLE_EXP: { os << (const struct ur_ipc_close_mem_handle_exp_params_t *)params; } break; + case UR_FUNCTION_IPC_GET_PHYS_MEM_HANDLE_EXP: { + os << (const struct ur_ipc_get_phys_mem_handle_exp_params_t *)params; + } break; + case UR_FUNCTION_IPC_PUT_PHYS_MEM_HANDLE_EXP: { + os << (const struct ur_ipc_put_phys_mem_handle_exp_params_t *)params; + } break; + case UR_FUNCTION_IPC_OPEN_PHYS_MEM_HANDLE_EXP: { + os << (const struct ur_ipc_open_phys_mem_handle_exp_params_t *)params; + } break; + case UR_FUNCTION_IPC_CLOSE_PHYS_MEM_HANDLE_EXP: { + os << (const struct ur_ipc_close_phys_mem_handle_exp_params_t *)params; + } break; case UR_FUNCTION_MEMORY_EXPORT_ALLOC_EXPORTABLE_MEMORY_EXP: { os << (const struct ur_memory_export_alloc_exportable_memory_exp_params_t *) params; diff --git a/unified-runtime/source/adapters/level_zero/ur_interface_loader.cpp b/unified-runtime/source/adapters/level_zero/ur_interface_loader.cpp index 877a4d199c44..07004ee97714 100644 --- a/unified-runtime/source/adapters/level_zero/ur_interface_loader.cpp +++ b/unified-runtime/source/adapters/level_zero/ur_interface_loader.cpp @@ -291,6 +291,12 @@ UR_APIEXPORT ur_result_t UR_APICALL urGetIPCExpProcAddrTable( pDdiTable->pfnPutMemHandleExp = ur::level_zero::urIPCPutMemHandleExp; pDdiTable->pfnOpenMemHandleExp = ur::level_zero::urIPCOpenMemHandleExp; pDdiTable->pfnCloseMemHandleExp = ur::level_zero::urIPCCloseMemHandleExp; + pDdiTable->pfnGetPhysMemHandleExp = ur::level_zero::urIPCGetPhysMemHandleExp; + pDdiTable->pfnPutPhysMemHandleExp = ur::level_zero::urIPCPutPhysMemHandleExp; + pDdiTable->pfnOpenPhysMemHandleExp = + ur::level_zero::urIPCOpenPhysMemHandleExp; + pDdiTable->pfnClosePhysMemHandleExp = + ur::level_zero::urIPCClosePhysMemHandleExp; return result; } diff --git a/unified-runtime/source/adapters/level_zero/ur_interface_loader.hpp b/unified-runtime/source/adapters/level_zero/ur_interface_loader.hpp index a4a9e9819a95..62904d267db5 100644 --- a/unified-runtime/source/adapters/level_zero/ur_interface_loader.hpp +++ b/unified-runtime/source/adapters/level_zero/ur_interface_loader.hpp @@ -629,6 +629,19 @@ ur_result_t urIPCOpenMemHandleExp(ur_context_handle_t hContext, void *pIPCMemHandleData, size_t ipcMemHandleDataSize, void **ppMem); ur_result_t urIPCCloseMemHandleExp(ur_context_handle_t hContext, void *pMem); +ur_result_t urIPCGetPhysMemHandleExp(ur_context_handle_t hContext, + ur_physical_mem_handle_t hPhysMem, + void **ppIPCPhysMemHandleData, + size_t *pIPCPhysMemHandleDataSizeRet); +ur_result_t urIPCPutPhysMemHandleExp(ur_context_handle_t hContext, + void *pIPCPhysMemHandleData); +ur_result_t urIPCOpenPhysMemHandleExp(ur_context_handle_t hContext, + ur_device_handle_t hDevice, + void *pIPCPhysMemHandleData, + size_t ipcPhysMemHandleDataSize, + ur_physical_mem_handle_t *phPhysMem); +ur_result_t urIPCClosePhysMemHandleExp(ur_context_handle_t hContext, + ur_physical_mem_handle_t hPhysMem); ur_result_t urMemoryExportAllocExportableMemoryExp( ur_context_handle_t hContext, ur_device_handle_t hDevice, size_t alignment, size_t size, ur_exp_external_mem_type_t handleTypeToExport, void **ppMem); diff --git a/unified-runtime/source/adapters/mock/ur_mockddi.cpp b/unified-runtime/source/adapters/mock/ur_mockddi.cpp index 00145e81d7af..4ce53b06524a 100644 --- a/unified-runtime/source/adapters/mock/ur_mockddi.cpp +++ b/unified-runtime/source/adapters/mock/ur_mockddi.cpp @@ -9661,6 +9661,198 @@ __urdlllocal ur_result_t UR_APICALL urIPCCloseMemHandleExp( return exceptionToResult(std::current_exception()); } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCGetPhysMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCGetPhysMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the physical memory object + ur_physical_mem_handle_t hPhysMem, + /// [out][optional] a pointer to the IPC physical memory handle data + void **ppIPCPhysMemHandleData, + /// [out][optional] size of the resulting IPC physical memory handle data + size_t *pIPCPhysMemHandleDataSizeRet) try { + ur_result_t result = UR_RESULT_SUCCESS; + + ur_ipc_get_phys_mem_handle_exp_params_t params = { + &hContext, &hPhysMem, &ppIPCPhysMemHandleData, + &pIPCPhysMemHandleDataSizeRet}; + + auto beforeCallback = reinterpret_cast( + mock::getCallbacks().get_before_callback("urIPCGetPhysMemHandleExp")); + if (beforeCallback) { + result = beforeCallback(¶ms); + if (result != UR_RESULT_SUCCESS) { + return result; + } + } + + auto replaceCallback = reinterpret_cast( + mock::getCallbacks().get_replace_callback("urIPCGetPhysMemHandleExp")); + if (replaceCallback) { + result = replaceCallback(¶ms); + } else { + + result = UR_RESULT_SUCCESS; + } + + if (result != UR_RESULT_SUCCESS) { + return result; + } + + auto afterCallback = reinterpret_cast( + mock::getCallbacks().get_after_callback("urIPCGetPhysMemHandleExp")); + if (afterCallback) { + return afterCallback(¶ms); + } + + return result; +} catch (...) { + return exceptionToResult(std::current_exception()); +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCPutPhysMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCPutPhysMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] a pointer to the IPC physical memory handle data + void *pIPCPhysMemHandleData) try { + ur_result_t result = UR_RESULT_SUCCESS; + + ur_ipc_put_phys_mem_handle_exp_params_t params = {&hContext, + &pIPCPhysMemHandleData}; + + auto beforeCallback = reinterpret_cast( + mock::getCallbacks().get_before_callback("urIPCPutPhysMemHandleExp")); + if (beforeCallback) { + result = beforeCallback(¶ms); + if (result != UR_RESULT_SUCCESS) { + return result; + } + } + + auto replaceCallback = reinterpret_cast( + mock::getCallbacks().get_replace_callback("urIPCPutPhysMemHandleExp")); + if (replaceCallback) { + result = replaceCallback(¶ms); + } else { + + result = UR_RESULT_SUCCESS; + } + + if (result != UR_RESULT_SUCCESS) { + return result; + } + + auto afterCallback = reinterpret_cast( + mock::getCallbacks().get_after_callback("urIPCPutPhysMemHandleExp")); + if (afterCallback) { + return afterCallback(¶ms); + } + + return result; +} catch (...) { + return exceptionToResult(std::current_exception()); +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCOpenPhysMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCOpenPhysMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object the physical memory was allocated on + ur_device_handle_t hDevice, + /// [in] the IPC physical memory handle data + void *pIPCPhysMemHandleData, + /// [in] size of the IPC physical memory handle data + size_t ipcPhysMemHandleDataSize, + /// [out] pointer to the physical memory handle + ur_physical_mem_handle_t *phPhysMem) try { + ur_result_t result = UR_RESULT_SUCCESS; + + ur_ipc_open_phys_mem_handle_exp_params_t params = { + &hContext, &hDevice, &pIPCPhysMemHandleData, &ipcPhysMemHandleDataSize, + &phPhysMem}; + + auto beforeCallback = reinterpret_cast( + mock::getCallbacks().get_before_callback("urIPCOpenPhysMemHandleExp")); + if (beforeCallback) { + result = beforeCallback(¶ms); + if (result != UR_RESULT_SUCCESS) { + return result; + } + } + + auto replaceCallback = reinterpret_cast( + mock::getCallbacks().get_replace_callback("urIPCOpenPhysMemHandleExp")); + if (replaceCallback) { + result = replaceCallback(¶ms); + } else { + + *phPhysMem = mock::createDummyHandle(); + result = UR_RESULT_SUCCESS; + } + + if (result != UR_RESULT_SUCCESS) { + return result; + } + + auto afterCallback = reinterpret_cast( + mock::getCallbacks().get_after_callback("urIPCOpenPhysMemHandleExp")); + if (afterCallback) { + return afterCallback(¶ms); + } + + return result; +} catch (...) { + return exceptionToResult(std::current_exception()); +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCClosePhysMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCClosePhysMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] physical memory handle opened through urIPCOpenPhysMemHandleExp + ur_physical_mem_handle_t hPhysMem) try { + ur_result_t result = UR_RESULT_SUCCESS; + + ur_ipc_close_phys_mem_handle_exp_params_t params = {&hContext, &hPhysMem}; + + auto beforeCallback = reinterpret_cast( + mock::getCallbacks().get_before_callback("urIPCClosePhysMemHandleExp")); + if (beforeCallback) { + result = beforeCallback(¶ms); + if (result != UR_RESULT_SUCCESS) { + return result; + } + } + + auto replaceCallback = reinterpret_cast( + mock::getCallbacks().get_replace_callback("urIPCClosePhysMemHandleExp")); + if (replaceCallback) { + result = replaceCallback(¶ms); + } else { + + result = UR_RESULT_SUCCESS; + } + + if (result != UR_RESULT_SUCCESS) { + return result; + } + + auto afterCallback = reinterpret_cast( + mock::getCallbacks().get_after_callback("urIPCClosePhysMemHandleExp")); + if (afterCallback) { + return afterCallback(¶ms); + } + + return result; +} catch (...) { + return exceptionToResult(std::current_exception()); +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Intercept function for urMemoryExportAllocExportableMemoryExp __urdlllocal ur_result_t UR_APICALL urMemoryExportAllocExportableMemoryExp( @@ -13492,6 +13684,14 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetIPCExpProcAddrTable( pDdiTable->pfnCloseMemHandleExp = driver::urIPCCloseMemHandleExp; + pDdiTable->pfnGetPhysMemHandleExp = driver::urIPCGetPhysMemHandleExp; + + pDdiTable->pfnPutPhysMemHandleExp = driver::urIPCPutPhysMemHandleExp; + + pDdiTable->pfnOpenPhysMemHandleExp = driver::urIPCOpenPhysMemHandleExp; + + pDdiTable->pfnClosePhysMemHandleExp = driver::urIPCClosePhysMemHandleExp; + return result; } catch (...) { return exceptionToResult(std::current_exception()); diff --git a/unified-runtime/source/adapters/opencl/common.hpp b/unified-runtime/source/adapters/opencl/common.hpp index aefea1f6c55a..fadab9d19992 100644 --- a/unified-runtime/source/adapters/opencl/common.hpp +++ b/unified-runtime/source/adapters/opencl/common.hpp @@ -224,24 +224,24 @@ CONSTFIX char GetKernelSubGroupInfoName[] = "clGetKernelSubGroupInfoKHR"; #undef CONSTFIX using clGetDeviceFunctionPointerINTEL_fn = CL_API_ENTRY - cl_int(CL_API_CALL *)(cl_device_id device, cl_program program, - const char *FuncName, cl_ulong *ret_ptr); +cl_int(CL_API_CALL *)(cl_device_id device, cl_program program, + const char *FuncName, cl_ulong *ret_ptr); using clGetDeviceGlobalVariablePointerINTEL_fn = CL_API_ENTRY - cl_int(CL_API_CALL *)(cl_device_id device, cl_program program, - const char *globalVariableName, - size_t *globalVariableSizeRet, - void **globalVariablePointerRet); +cl_int(CL_API_CALL *)(cl_device_id device, cl_program program, + const char *globalVariableName, + size_t *globalVariableSizeRet, + void **globalVariablePointerRet); using clEnqueueWriteGlobalVariableINTEL_fn = CL_API_ENTRY - cl_int(CL_API_CALL *)(cl_command_queue, cl_program, const char *, cl_bool, - size_t, size_t, const void *, cl_uint, - const cl_event *, cl_event *); +cl_int(CL_API_CALL *)(cl_command_queue, cl_program, const char *, cl_bool, + size_t, size_t, const void *, cl_uint, const cl_event *, + cl_event *); using clEnqueueReadGlobalVariableINTEL_fn = CL_API_ENTRY - cl_int(CL_API_CALL *)(cl_command_queue, cl_program, const char *, cl_bool, - size_t, size_t, void *, cl_uint, const cl_event *, - cl_event *); +cl_int(CL_API_CALL *)(cl_command_queue, cl_program, const char *, cl_bool, + size_t, size_t, void *, cl_uint, const cl_event *, + cl_event *); using clEnqueueReadHostPipeINTEL_fn = CL_API_ENTRY cl_int(CL_API_CALL *)(cl_command_queue queue, cl_program program, diff --git a/unified-runtime/source/loader/layers/tracing/ur_trcddi.cpp b/unified-runtime/source/loader/layers/tracing/ur_trcddi.cpp index 6d6ac7919f4d..e03069021909 100644 --- a/unified-runtime/source/loader/layers/tracing/ur_trcddi.cpp +++ b/unified-runtime/source/loader/layers/tracing/ur_trcddi.cpp @@ -8159,6 +8159,176 @@ __urdlllocal ur_result_t UR_APICALL urIPCCloseMemHandleExp( return result; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCGetPhysMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCGetPhysMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the physical memory object + ur_physical_mem_handle_t hPhysMem, + /// [out][optional] a pointer to the IPC physical memory handle data + void **ppIPCPhysMemHandleData, + /// [out][optional] size of the resulting IPC physical memory handle data + size_t *pIPCPhysMemHandleDataSizeRet) { + auto pfnGetPhysMemHandleExp = + getContext()->urDdiTable.IPCExp.pfnGetPhysMemHandleExp; + + if (nullptr == pfnGetPhysMemHandleExp) + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + + ur_ipc_get_phys_mem_handle_exp_params_t params = { + &hContext, &hPhysMem, &ppIPCPhysMemHandleData, + &pIPCPhysMemHandleDataSizeRet}; + uint64_t instance = + getContext()->notify_begin(UR_FUNCTION_IPC_GET_PHYS_MEM_HANDLE_EXP, + "urIPCGetPhysMemHandleExp", ¶ms); + + auto &logger = getContext()->logger; + UR_LOG_L(logger, INFO, " ---> urIPCGetPhysMemHandleExp\n"); + + ur_result_t result = pfnGetPhysMemHandleExp( + hContext, hPhysMem, ppIPCPhysMemHandleData, pIPCPhysMemHandleDataSizeRet); + + getContext()->notify_end(UR_FUNCTION_IPC_GET_PHYS_MEM_HANDLE_EXP, + "urIPCGetPhysMemHandleExp", ¶ms, &result, + instance); + + if (logger.getLevel() <= UR_LOGGER_LEVEL_INFO) { + std::ostringstream args_str; + ur::extras::printFunctionParams( + args_str, UR_FUNCTION_IPC_GET_PHYS_MEM_HANDLE_EXP, ¶ms); + UR_LOG_L(logger, INFO, " <--- urIPCGetPhysMemHandleExp({}) -> {};\n", + args_str.str(), result); + } + + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCPutPhysMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCPutPhysMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] a pointer to the IPC physical memory handle data + void *pIPCPhysMemHandleData) { + auto pfnPutPhysMemHandleExp = + getContext()->urDdiTable.IPCExp.pfnPutPhysMemHandleExp; + + if (nullptr == pfnPutPhysMemHandleExp) + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + + ur_ipc_put_phys_mem_handle_exp_params_t params = {&hContext, + &pIPCPhysMemHandleData}; + uint64_t instance = + getContext()->notify_begin(UR_FUNCTION_IPC_PUT_PHYS_MEM_HANDLE_EXP, + "urIPCPutPhysMemHandleExp", ¶ms); + + auto &logger = getContext()->logger; + UR_LOG_L(logger, INFO, " ---> urIPCPutPhysMemHandleExp\n"); + + ur_result_t result = pfnPutPhysMemHandleExp(hContext, pIPCPhysMemHandleData); + + getContext()->notify_end(UR_FUNCTION_IPC_PUT_PHYS_MEM_HANDLE_EXP, + "urIPCPutPhysMemHandleExp", ¶ms, &result, + instance); + + if (logger.getLevel() <= UR_LOGGER_LEVEL_INFO) { + std::ostringstream args_str; + ur::extras::printFunctionParams( + args_str, UR_FUNCTION_IPC_PUT_PHYS_MEM_HANDLE_EXP, ¶ms); + UR_LOG_L(logger, INFO, " <--- urIPCPutPhysMemHandleExp({}) -> {};\n", + args_str.str(), result); + } + + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCOpenPhysMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCOpenPhysMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object the physical memory was allocated on + ur_device_handle_t hDevice, + /// [in] the IPC physical memory handle data + void *pIPCPhysMemHandleData, + /// [in] size of the IPC physical memory handle data + size_t ipcPhysMemHandleDataSize, + /// [out] pointer to the physical memory handle + ur_physical_mem_handle_t *phPhysMem) { + auto pfnOpenPhysMemHandleExp = + getContext()->urDdiTable.IPCExp.pfnOpenPhysMemHandleExp; + + if (nullptr == pfnOpenPhysMemHandleExp) + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + + ur_ipc_open_phys_mem_handle_exp_params_t params = { + &hContext, &hDevice, &pIPCPhysMemHandleData, &ipcPhysMemHandleDataSize, + &phPhysMem}; + uint64_t instance = + getContext()->notify_begin(UR_FUNCTION_IPC_OPEN_PHYS_MEM_HANDLE_EXP, + "urIPCOpenPhysMemHandleExp", ¶ms); + + auto &logger = getContext()->logger; + UR_LOG_L(logger, INFO, " ---> urIPCOpenPhysMemHandleExp\n"); + + ur_result_t result = + pfnOpenPhysMemHandleExp(hContext, hDevice, pIPCPhysMemHandleData, + ipcPhysMemHandleDataSize, phPhysMem); + + getContext()->notify_end(UR_FUNCTION_IPC_OPEN_PHYS_MEM_HANDLE_EXP, + "urIPCOpenPhysMemHandleExp", ¶ms, &result, + instance); + + if (logger.getLevel() <= UR_LOGGER_LEVEL_INFO) { + std::ostringstream args_str; + ur::extras::printFunctionParams( + args_str, UR_FUNCTION_IPC_OPEN_PHYS_MEM_HANDLE_EXP, ¶ms); + UR_LOG_L(logger, INFO, " <--- urIPCOpenPhysMemHandleExp({}) -> {};\n", + args_str.str(), result); + } + + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCClosePhysMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCClosePhysMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] physical memory handle opened through urIPCOpenPhysMemHandleExp + ur_physical_mem_handle_t hPhysMem) { + auto pfnClosePhysMemHandleExp = + getContext()->urDdiTable.IPCExp.pfnClosePhysMemHandleExp; + + if (nullptr == pfnClosePhysMemHandleExp) + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + + ur_ipc_close_phys_mem_handle_exp_params_t params = {&hContext, &hPhysMem}; + uint64_t instance = + getContext()->notify_begin(UR_FUNCTION_IPC_CLOSE_PHYS_MEM_HANDLE_EXP, + "urIPCClosePhysMemHandleExp", ¶ms); + + auto &logger = getContext()->logger; + UR_LOG_L(logger, INFO, " ---> urIPCClosePhysMemHandleExp\n"); + + ur_result_t result = pfnClosePhysMemHandleExp(hContext, hPhysMem); + + getContext()->notify_end(UR_FUNCTION_IPC_CLOSE_PHYS_MEM_HANDLE_EXP, + "urIPCClosePhysMemHandleExp", ¶ms, &result, + instance); + + if (logger.getLevel() <= UR_LOGGER_LEVEL_INFO) { + std::ostringstream args_str; + ur::extras::printFunctionParams( + args_str, UR_FUNCTION_IPC_CLOSE_PHYS_MEM_HANDLE_EXP, ¶ms); + UR_LOG_L(logger, INFO, " <--- urIPCClosePhysMemHandleExp({}) -> {};\n", + args_str.str(), result); + } + + return result; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Intercept function for urMemoryExportAllocExportableMemoryExp __urdlllocal ur_result_t UR_APICALL urMemoryExportAllocExportableMemoryExp( @@ -11681,6 +11851,22 @@ __urdlllocal ur_result_t UR_APICALL urGetIPCExpProcAddrTable( dditable.pfnCloseMemHandleExp = pDdiTable->pfnCloseMemHandleExp; pDdiTable->pfnCloseMemHandleExp = ur_tracing_layer::urIPCCloseMemHandleExp; + dditable.pfnGetPhysMemHandleExp = pDdiTable->pfnGetPhysMemHandleExp; + pDdiTable->pfnGetPhysMemHandleExp = + ur_tracing_layer::urIPCGetPhysMemHandleExp; + + dditable.pfnPutPhysMemHandleExp = pDdiTable->pfnPutPhysMemHandleExp; + pDdiTable->pfnPutPhysMemHandleExp = + ur_tracing_layer::urIPCPutPhysMemHandleExp; + + dditable.pfnOpenPhysMemHandleExp = pDdiTable->pfnOpenPhysMemHandleExp; + pDdiTable->pfnOpenPhysMemHandleExp = + ur_tracing_layer::urIPCOpenPhysMemHandleExp; + + dditable.pfnClosePhysMemHandleExp = pDdiTable->pfnClosePhysMemHandleExp; + pDdiTable->pfnClosePhysMemHandleExp = + ur_tracing_layer::urIPCClosePhysMemHandleExp; + return result; } /////////////////////////////////////////////////////////////////////////////// diff --git a/unified-runtime/source/loader/layers/validation/ur_valddi.cpp b/unified-runtime/source/loader/layers/validation/ur_valddi.cpp index 7b9a2f5b0982..32ace25e959a 100644 --- a/unified-runtime/source/loader/layers/validation/ur_valddi.cpp +++ b/unified-runtime/source/loader/layers/validation/ur_valddi.cpp @@ -9003,6 +9003,174 @@ __urdlllocal ur_result_t UR_APICALL urIPCCloseMemHandleExp( return result; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCGetPhysMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCGetPhysMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the physical memory object + ur_physical_mem_handle_t hPhysMem, + /// [out][optional] a pointer to the IPC physical memory handle data + void **ppIPCPhysMemHandleData, + /// [out][optional] size of the resulting IPC physical memory handle data + size_t *pIPCPhysMemHandleDataSizeRet) { + auto pfnGetPhysMemHandleExp = + getContext()->urDdiTable.IPCExp.pfnGetPhysMemHandleExp; + + if (nullptr == pfnGetPhysMemHandleExp) { + return UR_RESULT_ERROR_UNINITIALIZED; + } + + if (getContext()->enableParameterValidation) { + if (NULL == ppIPCPhysMemHandleData) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + + if (NULL == pIPCPhysMemHandleDataSizeRet) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + + if (NULL == hContext) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + + if (NULL == hPhysMem) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + } + + if (getContext()->enableLifetimeValidation && + !getContext()->refCountContext->isReferenceValid(hContext)) { + URLOG_CTX_INVALID_REFERENCE(hContext); + } + + if (getContext()->enableLifetimeValidation && + !getContext()->refCountContext->isReferenceValid(hPhysMem)) { + URLOG_CTX_INVALID_REFERENCE(hPhysMem); + } + + ur_result_t result = pfnGetPhysMemHandleExp( + hContext, hPhysMem, ppIPCPhysMemHandleData, pIPCPhysMemHandleDataSizeRet); + + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCPutPhysMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCPutPhysMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] a pointer to the IPC physical memory handle data + void *pIPCPhysMemHandleData) { + auto pfnPutPhysMemHandleExp = + getContext()->urDdiTable.IPCExp.pfnPutPhysMemHandleExp; + + if (nullptr == pfnPutPhysMemHandleExp) { + return UR_RESULT_ERROR_UNINITIALIZED; + } + + if (getContext()->enableParameterValidation) { + if (NULL == pIPCPhysMemHandleData) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + + if (NULL == hContext) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + } + + if (getContext()->enableLifetimeValidation && + !getContext()->refCountContext->isReferenceValid(hContext)) { + URLOG_CTX_INVALID_REFERENCE(hContext); + } + + ur_result_t result = pfnPutPhysMemHandleExp(hContext, pIPCPhysMemHandleData); + + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCOpenPhysMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCOpenPhysMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object the physical memory was allocated on + ur_device_handle_t hDevice, + /// [in] the IPC physical memory handle data + void *pIPCPhysMemHandleData, + /// [in] size of the IPC physical memory handle data + size_t ipcPhysMemHandleDataSize, + /// [out] pointer to the physical memory handle + ur_physical_mem_handle_t *phPhysMem) { + auto pfnOpenPhysMemHandleExp = + getContext()->urDdiTable.IPCExp.pfnOpenPhysMemHandleExp; + + if (nullptr == pfnOpenPhysMemHandleExp) { + return UR_RESULT_ERROR_UNINITIALIZED; + } + + if (getContext()->enableParameterValidation) { + if (NULL == phPhysMem) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + + if (NULL == pIPCPhysMemHandleData) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + + if (NULL == hContext) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + + if (NULL == hDevice) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + } + + if (getContext()->enableLifetimeValidation && + !getContext()->refCountContext->isReferenceValid(hContext)) { + URLOG_CTX_INVALID_REFERENCE(hContext); + } + + if (getContext()->enableLifetimeValidation && + !getContext()->refCountContext->isReferenceValid(hDevice)) { + URLOG_CTX_INVALID_REFERENCE(hDevice); + } + + ur_result_t result = + pfnOpenPhysMemHandleExp(hContext, hDevice, pIPCPhysMemHandleData, + ipcPhysMemHandleDataSize, phPhysMem); + + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCClosePhysMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCClosePhysMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] physical memory handle opened through urIPCOpenPhysMemHandleExp + ur_physical_mem_handle_t hPhysMem) { + auto pfnClosePhysMemHandleExp = + getContext()->urDdiTable.IPCExp.pfnClosePhysMemHandleExp; + + if (nullptr == pfnClosePhysMemHandleExp) { + return UR_RESULT_ERROR_UNINITIALIZED; + } + + if (getContext()->enableParameterValidation) { + if (NULL == hContext) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + + if (NULL == hPhysMem) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + } + + if (getContext()->enableLifetimeValidation && + !getContext()->refCountContext->isReferenceValid(hContext)) { + URLOG_CTX_INVALID_REFERENCE(hContext); + } + + if (getContext()->enableLifetimeValidation && + !getContext()->refCountContext->isReferenceValid(hPhysMem)) { + URLOG_CTX_INVALID_REFERENCE(hPhysMem); + } + + ur_result_t result = pfnClosePhysMemHandleExp(hContext, hPhysMem); + + return result; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Intercept function for urMemoryExportAllocExportableMemoryExp __urdlllocal ur_result_t UR_APICALL urMemoryExportAllocExportableMemoryExp( @@ -12465,6 +12633,22 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetIPCExpProcAddrTable( dditable.pfnCloseMemHandleExp = pDdiTable->pfnCloseMemHandleExp; pDdiTable->pfnCloseMemHandleExp = ur_validation_layer::urIPCCloseMemHandleExp; + dditable.pfnGetPhysMemHandleExp = pDdiTable->pfnGetPhysMemHandleExp; + pDdiTable->pfnGetPhysMemHandleExp = + ur_validation_layer::urIPCGetPhysMemHandleExp; + + dditable.pfnPutPhysMemHandleExp = pDdiTable->pfnPutPhysMemHandleExp; + pDdiTable->pfnPutPhysMemHandleExp = + ur_validation_layer::urIPCPutPhysMemHandleExp; + + dditable.pfnOpenPhysMemHandleExp = pDdiTable->pfnOpenPhysMemHandleExp; + pDdiTable->pfnOpenPhysMemHandleExp = + ur_validation_layer::urIPCOpenPhysMemHandleExp; + + dditable.pfnClosePhysMemHandleExp = pDdiTable->pfnClosePhysMemHandleExp; + pDdiTable->pfnClosePhysMemHandleExp = + ur_validation_layer::urIPCClosePhysMemHandleExp; + return result; } diff --git a/unified-runtime/source/loader/loader.def.in b/unified-runtime/source/loader/loader.def.in index 8faa938b745f..c31f6c33bbe6 100644 --- a/unified-runtime/source/loader/loader.def.in +++ b/unified-runtime/source/loader/loader.def.in @@ -146,9 +146,13 @@ EXPORTS urGraphInstantiateGraphExp urGraphIsEmptyExp urIPCCloseMemHandleExp + urIPCClosePhysMemHandleExp urIPCGetMemHandleExp + urIPCGetPhysMemHandleExp urIPCOpenMemHandleExp + urIPCOpenPhysMemHandleExp urIPCPutMemHandleExp + urIPCPutPhysMemHandleExp urKernelCreate urKernelCreateWithNativeHandle urKernelGetGroupInfo @@ -395,9 +399,13 @@ EXPORTS urPrintImageFormat urPrintImageInfo urPrintIpcCloseMemHandleExpParams + urPrintIpcClosePhysMemHandleExpParams urPrintIpcGetMemHandleExpParams + urPrintIpcGetPhysMemHandleExpParams urPrintIpcOpenMemHandleExpParams + urPrintIpcOpenPhysMemHandleExpParams urPrintIpcPutMemHandleExpParams + urPrintIpcPutPhysMemHandleExpParams urPrintKernelArgLocalProperties urPrintKernelArgMemObjProperties urPrintKernelArgPointerProperties diff --git a/unified-runtime/source/loader/loader.map.in b/unified-runtime/source/loader/loader.map.in index b7f9f3428a5f..b04224a96d86 100644 --- a/unified-runtime/source/loader/loader.map.in +++ b/unified-runtime/source/loader/loader.map.in @@ -146,9 +146,13 @@ urGraphInstantiateGraphExp; urGraphIsEmptyExp; urIPCCloseMemHandleExp; + urIPCClosePhysMemHandleExp; urIPCGetMemHandleExp; + urIPCGetPhysMemHandleExp; urIPCOpenMemHandleExp; + urIPCOpenPhysMemHandleExp; urIPCPutMemHandleExp; + urIPCPutPhysMemHandleExp; urKernelCreate; urKernelCreateWithNativeHandle; urKernelGetGroupInfo; @@ -395,9 +399,13 @@ urPrintImageFormat; urPrintImageInfo; urPrintIpcCloseMemHandleExpParams; + urPrintIpcClosePhysMemHandleExpParams; urPrintIpcGetMemHandleExpParams; + urPrintIpcGetPhysMemHandleExpParams; urPrintIpcOpenMemHandleExpParams; + urPrintIpcOpenPhysMemHandleExpParams; urPrintIpcPutMemHandleExpParams; + urPrintIpcPutPhysMemHandleExpParams; urPrintKernelArgLocalProperties; urPrintKernelArgMemObjProperties; urPrintKernelArgPointerProperties; diff --git a/unified-runtime/source/loader/ur_ldrddi.cpp b/unified-runtime/source/loader/ur_ldrddi.cpp index 7ac726f9c9cb..32a9307db62d 100644 --- a/unified-runtime/source/loader/ur_ldrddi.cpp +++ b/unified-runtime/source/loader/ur_ldrddi.cpp @@ -4643,6 +4643,90 @@ __urdlllocal ur_result_t UR_APICALL urIPCCloseMemHandleExp( return pfnCloseMemHandleExp(hContext, pMem); } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCGetPhysMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCGetPhysMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the physical memory object + ur_physical_mem_handle_t hPhysMem, + /// [out][optional] a pointer to the IPC physical memory handle data + void **ppIPCPhysMemHandleData, + /// [out][optional] size of the resulting IPC physical memory handle data + size_t *pIPCPhysMemHandleDataSizeRet) { + + auto *dditable = *reinterpret_cast(hContext); + + auto *pfnGetPhysMemHandleExp = dditable->IPCExp.pfnGetPhysMemHandleExp; + if (nullptr == pfnGetPhysMemHandleExp) + return UR_RESULT_ERROR_UNINITIALIZED; + + // forward to device-platform + return pfnGetPhysMemHandleExp(hContext, hPhysMem, ppIPCPhysMemHandleData, + pIPCPhysMemHandleDataSizeRet); +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCPutPhysMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCPutPhysMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] a pointer to the IPC physical memory handle data + void *pIPCPhysMemHandleData) { + + auto *dditable = *reinterpret_cast(hContext); + + auto *pfnPutPhysMemHandleExp = dditable->IPCExp.pfnPutPhysMemHandleExp; + if (nullptr == pfnPutPhysMemHandleExp) + return UR_RESULT_ERROR_UNINITIALIZED; + + // forward to device-platform + return pfnPutPhysMemHandleExp(hContext, pIPCPhysMemHandleData); +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCOpenPhysMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCOpenPhysMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object the physical memory was allocated on + ur_device_handle_t hDevice, + /// [in] the IPC physical memory handle data + void *pIPCPhysMemHandleData, + /// [in] size of the IPC physical memory handle data + size_t ipcPhysMemHandleDataSize, + /// [out] pointer to the physical memory handle + ur_physical_mem_handle_t *phPhysMem) { + + auto *dditable = *reinterpret_cast(hContext); + + auto *pfnOpenPhysMemHandleExp = dditable->IPCExp.pfnOpenPhysMemHandleExp; + if (nullptr == pfnOpenPhysMemHandleExp) + return UR_RESULT_ERROR_UNINITIALIZED; + + // forward to device-platform + return pfnOpenPhysMemHandleExp(hContext, hDevice, pIPCPhysMemHandleData, + ipcPhysMemHandleDataSize, phPhysMem); +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCClosePhysMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCClosePhysMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] physical memory handle opened through urIPCOpenPhysMemHandleExp + ur_physical_mem_handle_t hPhysMem) { + + auto *dditable = *reinterpret_cast(hContext); + + auto *pfnClosePhysMemHandleExp = dditable->IPCExp.pfnClosePhysMemHandleExp; + if (nullptr == pfnClosePhysMemHandleExp) + return UR_RESULT_ERROR_UNINITIALIZED; + + // forward to device-platform + return pfnClosePhysMemHandleExp(hContext, hPhysMem); +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Intercept function for urMemoryExportAllocExportableMemoryExp __urdlllocal ur_result_t UR_APICALL urMemoryExportAllocExportableMemoryExp( @@ -6874,6 +6958,11 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetIPCExpProcAddrTable( pDdiTable->pfnPutMemHandleExp = ur_loader::urIPCPutMemHandleExp; pDdiTable->pfnOpenMemHandleExp = ur_loader::urIPCOpenMemHandleExp; pDdiTable->pfnCloseMemHandleExp = ur_loader::urIPCCloseMemHandleExp; + pDdiTable->pfnGetPhysMemHandleExp = ur_loader::urIPCGetPhysMemHandleExp; + pDdiTable->pfnPutPhysMemHandleExp = ur_loader::urIPCPutPhysMemHandleExp; + pDdiTable->pfnOpenPhysMemHandleExp = ur_loader::urIPCOpenPhysMemHandleExp; + pDdiTable->pfnClosePhysMemHandleExp = + ur_loader::urIPCClosePhysMemHandleExp; } else { // return pointers directly to platform's DDIs *pDdiTable = ur_loader::getContext()->platforms.front().dditable.IPCExp; diff --git a/unified-runtime/source/loader/ur_libapi.cpp b/unified-runtime/source/loader/ur_libapi.cpp index 51f6ca8ee556..25ef19728492 100644 --- a/unified-runtime/source/loader/ur_libapi.cpp +++ b/unified-runtime/source/loader/ur_libapi.cpp @@ -8640,6 +8640,146 @@ ur_result_t UR_APICALL urIPCCloseMemHandleExp( return exceptionToResult(std::current_exception()); } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Gets an inter-process handle for a physical memory object +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hContext` +/// + `NULL == hPhysMem` +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == ppIPCPhysMemHandleData` +/// + `NULL == pIPCPhysMemHandleDataSizeRet` +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +ur_result_t UR_APICALL urIPCGetPhysMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the physical memory object + ur_physical_mem_handle_t hPhysMem, + /// [out][optional] a pointer to the IPC physical memory handle data + void **ppIPCPhysMemHandleData, + /// [out][optional] size of the resulting IPC physical memory handle data + size_t *pIPCPhysMemHandleDataSizeRet) try { + auto pfnGetPhysMemHandleExp = + ur_lib::getContext()->urDdiTable.IPCExp.pfnGetPhysMemHandleExp; + if (nullptr == pfnGetPhysMemHandleExp) + return UR_RESULT_ERROR_UNINITIALIZED; + + return pfnGetPhysMemHandleExp(hContext, hPhysMem, ppIPCPhysMemHandleData, + pIPCPhysMemHandleDataSizeRet); +} catch (...) { + return exceptionToResult(std::current_exception()); +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Releases an inter-process physical memory handle +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hContext` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pIPCPhysMemHandleData` +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +ur_result_t UR_APICALL urIPCPutPhysMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] a pointer to the IPC physical memory handle data + void *pIPCPhysMemHandleData) try { + auto pfnPutPhysMemHandleExp = + ur_lib::getContext()->urDdiTable.IPCExp.pfnPutPhysMemHandleExp; + if (nullptr == pfnPutPhysMemHandleExp) + return UR_RESULT_ERROR_UNINITIALIZED; + + return pfnPutPhysMemHandleExp(hContext, pIPCPhysMemHandleData); +} catch (...) { + return exceptionToResult(std::current_exception()); +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Opens an inter-process physical memory handle to get the +/// corresponding +/// physical memory object +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hContext` +/// + `NULL == hDevice` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == phPhysMem` +/// + `NULL == pIPCPhysMemHandleData` +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +/// - ::UR_RESULT_ERROR_INVALID_VALUE +/// + ipcPhysMemHandleDataSize is not the same as the size of IPC +/// physical memory handle data +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +ur_result_t UR_APICALL urIPCOpenPhysMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object the physical memory was allocated on + ur_device_handle_t hDevice, + /// [in] the IPC physical memory handle data + void *pIPCPhysMemHandleData, + /// [in] size of the IPC physical memory handle data + size_t ipcPhysMemHandleDataSize, + /// [out] pointer to the physical memory handle + ur_physical_mem_handle_t *phPhysMem) try { + auto pfnOpenPhysMemHandleExp = + ur_lib::getContext()->urDdiTable.IPCExp.pfnOpenPhysMemHandleExp; + if (nullptr == pfnOpenPhysMemHandleExp) + return UR_RESULT_ERROR_UNINITIALIZED; + + return pfnOpenPhysMemHandleExp(hContext, hDevice, pIPCPhysMemHandleData, + ipcPhysMemHandleDataSize, phPhysMem); +} catch (...) { + return exceptionToResult(std::current_exception()); +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Closes an inter-process physical memory handle +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hContext` +/// + `NULL == hPhysMem` +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +ur_result_t UR_APICALL urIPCClosePhysMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] physical memory handle opened through urIPCOpenPhysMemHandleExp + ur_physical_mem_handle_t hPhysMem) try { + auto pfnClosePhysMemHandleExp = + ur_lib::getContext()->urDdiTable.IPCExp.pfnClosePhysMemHandleExp; + if (nullptr == pfnClosePhysMemHandleExp) + return UR_RESULT_ERROR_UNINITIALIZED; + + return pfnClosePhysMemHandleExp(hContext, hPhysMem); +} catch (...) { + return exceptionToResult(std::current_exception()); +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Allocate an exportable memory region and return a pointer to that /// allocation. diff --git a/unified-runtime/source/loader/ur_print.cpp b/unified-runtime/source/loader/ur_print.cpp index f500efa52e86..98d2ebc83a80 100644 --- a/unified-runtime/source/loader/ur_print.cpp +++ b/unified-runtime/source/loader/ur_print.cpp @@ -2195,6 +2195,38 @@ ur_result_t urPrintIpcCloseMemHandleExpParams( return str_copy(&ss, buffer, buff_size, out_size); } +ur_result_t urPrintIpcGetPhysMemHandleExpParams( + const struct ur_ipc_get_phys_mem_handle_exp_params_t *params, char *buffer, + const size_t buff_size, size_t *out_size) { + std::stringstream ss; + ss << params; + return str_copy(&ss, buffer, buff_size, out_size); +} + +ur_result_t urPrintIpcPutPhysMemHandleExpParams( + const struct ur_ipc_put_phys_mem_handle_exp_params_t *params, char *buffer, + const size_t buff_size, size_t *out_size) { + std::stringstream ss; + ss << params; + return str_copy(&ss, buffer, buff_size, out_size); +} + +ur_result_t urPrintIpcOpenPhysMemHandleExpParams( + const struct ur_ipc_open_phys_mem_handle_exp_params_t *params, char *buffer, + const size_t buff_size, size_t *out_size) { + std::stringstream ss; + ss << params; + return str_copy(&ss, buffer, buff_size, out_size); +} + +ur_result_t urPrintIpcClosePhysMemHandleExpParams( + const struct ur_ipc_close_phys_mem_handle_exp_params_t *params, + char *buffer, const size_t buff_size, size_t *out_size) { + std::stringstream ss; + ss << params; + return str_copy(&ss, buffer, buff_size, out_size); +} + ur_result_t urPrintKernelCreateParams(const struct ur_kernel_create_params_t *params, char *buffer, const size_t buff_size, diff --git a/unified-runtime/source/ur_api.cpp b/unified-runtime/source/ur_api.cpp index 81dcda1a4fcb..01640b1d7c3b 100644 --- a/unified-runtime/source/ur_api.cpp +++ b/unified-runtime/source/ur_api.cpp @@ -7522,6 +7522,120 @@ ur_result_t UR_APICALL urIPCCloseMemHandleExp( return result; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Gets an inter-process handle for a physical memory object +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hContext` +/// + `NULL == hPhysMem` +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == ppIPCPhysMemHandleData` +/// + `NULL == pIPCPhysMemHandleDataSizeRet` +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +ur_result_t UR_APICALL urIPCGetPhysMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the physical memory object + ur_physical_mem_handle_t hPhysMem, + /// [out][optional] a pointer to the IPC physical memory handle data + void **ppIPCPhysMemHandleData, + /// [out][optional] size of the resulting IPC physical memory handle data + size_t *pIPCPhysMemHandleDataSizeRet) { + ur_result_t result = UR_RESULT_SUCCESS; + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Releases an inter-process physical memory handle +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hContext` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pIPCPhysMemHandleData` +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +ur_result_t UR_APICALL urIPCPutPhysMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] a pointer to the IPC physical memory handle data + void *pIPCPhysMemHandleData) { + ur_result_t result = UR_RESULT_SUCCESS; + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Opens an inter-process physical memory handle to get the +/// corresponding +/// physical memory object +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hContext` +/// + `NULL == hDevice` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == phPhysMem` +/// + `NULL == pIPCPhysMemHandleData` +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +/// - ::UR_RESULT_ERROR_INVALID_VALUE +/// + ipcPhysMemHandleDataSize is not the same as the size of IPC +/// physical memory handle data +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +ur_result_t UR_APICALL urIPCOpenPhysMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object the physical memory was allocated on + ur_device_handle_t hDevice, + /// [in] the IPC physical memory handle data + void *pIPCPhysMemHandleData, + /// [in] size of the IPC physical memory handle data + size_t ipcPhysMemHandleDataSize, + /// [out] pointer to the physical memory handle + ur_physical_mem_handle_t *phPhysMem) { + ur_result_t result = UR_RESULT_SUCCESS; + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Closes an inter-process physical memory handle +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hContext` +/// + `NULL == hPhysMem` +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +ur_result_t UR_APICALL urIPCClosePhysMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] physical memory handle opened through urIPCOpenPhysMemHandleExp + ur_physical_mem_handle_t hPhysMem) { + ur_result_t result = UR_RESULT_SUCCESS; + return result; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Allocate an exportable memory region and return a pointer to that /// allocation. From 1c133b8646439eb15c93391172a7716c2c556b09 Mon Sep 17 00:00:00 2001 From: Lukasz Dorau Date: Wed, 20 May 2026 14:42:32 +0000 Subject: [PATCH 04/17] [SYCL][UR] Implement IPC API for physical_mem objects Add ipc::memory and ipc::physical_memory namespaces as the non-deprecated replacements for the old ipc_memory namespace, plus new IPC support for physical_mem objects (sycl_ext_oneapi_virtual_mem). SYCL layer changes: - ipc_memory.hpp: add ipc::handle generic handle type with proper forward-declaration pattern; add ipc::memory and ipc::physical_memory namespaces with get/put/open/close free functions; deprecate the old ipc_memory namespace - ipc_memory.cpp: implement ipc::memory::{get,put,close}, ipc::physical_memory::{get,put,open}, and detail::openIPCPhysMemHandle - physical_mem.hpp: add enable_ipc property struct and is_property_key_of specialization; add templated constructors that extract the property; add ext_oneapi_ipc_enabled() query method - physical_mem_impl.hpp: add bool EnableIpc parameter and MEnabledIpc member; add second constructor for opening from an existing UR handle (used by ipc::physical_memory::open) - physical_mem.cpp: propagate EnableIpc through to impl; add ext_oneapi_ipc_enabled(); check ext_oneapi_ipc_physical_memory aspect - aspects.def: add ext_oneapi_ipc_physical_memory (aspect 96) - property.hpp: add PropKind::EnableIpc = 50, bump PropKindSize to 51 UR adapter changes: - All adapter memory.cpp files: add empty stub implementations returning UR_RESULT_ERROR_UNSUPPORTED_FEATURE for the 4 new IPC functions - Non-level_zero adapter ur_interface_loader.cpp files: wire stubs into the IPC DDI table (after pfnCloseMemHandleExp) Signed-off-by: Lukasz Dorau --- .../oneapi/experimental/detail/ipc_common.hpp | 14 +++ .../ext/oneapi/experimental/ipc_memory.hpp | 61 ++++++++- .../sycl/ext/oneapi/properties/property.hpp | 3 +- .../ext/oneapi/virtual_mem/physical_mem.hpp | 52 +++++++- sycl/include/sycl/info/aspects.def | 1 + sycl/source/detail/physical_mem_impl.hpp | 22 +++- sycl/source/ipc_memory.cpp | 116 ++++++++++++++++++ sycl/source/physical_mem.cpp | 12 +- .../source/adapters/cuda/memory.cpp | 21 ++++ .../adapters/cuda/ur_interface_loader.cpp | 4 + .../source/adapters/hip/memory.cpp | 21 ++++ .../adapters/hip/ur_interface_loader.cpp | 4 + .../source/adapters/level_zero/memory.cpp | 21 ++++ .../source/adapters/level_zero/v2/memory.cpp | 21 ++++ .../source/adapters/native_cpu/memory.cpp | 21 ++++ .../native_cpu/ur_interface_loader.cpp | 4 + .../source/adapters/offload/memory.cpp | 21 ++++ .../adapters/offload/ur_interface_loader.cpp | 4 + .../source/adapters/opencl/memory.cpp | 21 ++++ .../adapters/opencl/ur_interface_loader.cpp | 4 + 20 files changed, 436 insertions(+), 12 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 a6cf6d481b8c..ad1c691baa3d 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/detail/ipc_common.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/detail/ipc_common.hpp @@ -24,6 +24,10 @@ namespace sycl { inline namespace _V1 { +namespace ext::oneapi::experimental { +class physical_mem; +} // namespace ext::oneapi::experimental + namespace ext::oneapi::experimental::ipc { struct handle; } @@ -33,6 +37,12 @@ __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::ipc::physical_memory { +__SYCL_EXPORT handle +get(const ext::oneapi::experimental::physical_mem &PhysMem); +__SYCL_EXPORT void put(handle &HandleData, const sycl::context &Ctx); +} // namespace ext::oneapi::experimental::ipc::physical_memory + namespace ext::oneapi::experimental::ipc { using handle_data_t = std::vector; @@ -59,6 +69,10 @@ 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(const ext::oneapi::experimental::physical_mem &PhysMem); + friend __SYCL_EXPORT void physical_memory::put(handle &HandleData, + const sycl::context &Ctx); }; } // namespace ext::oneapi::experimental::ipc diff --git a/sycl/include/sycl/ext/oneapi/experimental/ipc_memory.hpp b/sycl/include/sycl/ext/oneapi/experimental/ipc_memory.hpp index 1f1350942214..05a8cc94e93a 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/ipc_memory.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/ipc_memory.hpp @@ -14,6 +14,7 @@ #include #include #include +#include #include #include "detail/ipc_common.hpp" @@ -32,7 +33,11 @@ __SYCL_EXPORT void *openIPCMemHandle(const std::byte *HandleData, size_t HandleDataSize, const sycl::context &Ctx, const sycl::device &Dev); -} + +__SYCL_EXPORT ext::oneapi::experimental::physical_mem +openIPCPhysMemHandle(const std::byte *HandleData, size_t HandleDataSize, + const sycl::context &Ctx, const sycl::device &Dev); +} // namespace detail namespace ext::oneapi::experimental::ipc::memory { @@ -97,6 +102,60 @@ inline void close(void *Ptr) { } } // namespace ext::oneapi::experimental::ipc::memory +namespace ext::oneapi::experimental::ipc::physical_memory { + +__SYCL_EXPORT ipc::handle +get(const ext::oneapi::experimental::physical_mem &PhysMem); + +__SYCL_EXPORT void put(ipc::handle &HandleData, const sycl::context &Ctx); + +inline void put(ipc::handle &HandleData) { + sycl::device Dev; + sycl::context Ctx = Dev.get_platform().khr_get_default_context(); + ipc::physical_memory::put(HandleData, Ctx); +} + +__SYCL_EXPORT ext::oneapi::experimental::physical_mem +open(const ipc::handle_data_t &HandleData, const sycl::context &Ctx, + const sycl::device &Dev); + +inline ext::oneapi::experimental::physical_mem +open(const ipc::handle_data_t &HandleData, const sycl::device &Dev) { + sycl::context Ctx = Dev.get_platform().khr_get_default_context(); + return ipc::physical_memory::open(HandleData, Ctx, Dev); +} + +inline ext::oneapi::experimental::physical_mem +open(const ipc::handle_data_t &HandleData) { + sycl::device Dev; + sycl::context Ctx = Dev.get_platform().khr_get_default_context(); + return ipc::physical_memory::open(HandleData, Ctx, Dev); +} + +#if __cpp_lib_span +inline ext::oneapi::experimental::physical_mem +open(const ipc::handle_data_view_t &HandleDataView, const sycl::context &Ctx, + const sycl::device &Dev) { + ipc::handle_data_t HandleData{HandleDataView.begin(), HandleDataView.end()}; + return ipc::physical_memory::open(HandleData, Ctx, Dev); +} + +inline ext::oneapi::experimental::physical_mem +open(const ipc::handle_data_view_t &HandleDataView, const sycl::device &Dev) { + sycl::context Ctx = Dev.get_platform().khr_get_default_context(); + return ipc::physical_memory::open(HandleDataView, Ctx, Dev); +} + +inline ext::oneapi::experimental::physical_mem +open(const ipc::handle_data_view_t &HandleDataView) { + sycl::device Dev; + sycl::context Ctx = Dev.get_platform().khr_get_default_context(); + return ipc::physical_memory::open(HandleDataView, Ctx, Dev); +} +#endif + +} // namespace ext::oneapi::experimental::ipc::physical_memory + namespace ext::oneapi::experimental { namespace __SYCL_DEPRECATED("The ipc_memory namespace is deprecated. Use the " "ipc::memory namespace instead.") ipc_memory { diff --git a/sycl/include/sycl/ext/oneapi/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp index d73fdd6e5166..2d39e8684519 100644 --- a/sycl/include/sycl/ext/oneapi/properties/property.hpp +++ b/sycl/include/sycl/ext/oneapi/properties/property.hpp @@ -192,8 +192,9 @@ enum PropKind : uint32_t { MaximumSize = 47, ZeroInit = 48, FastLink = 49, + EnableIpc = 50, // PropKindSize must always be the last value. - PropKindSize = 50, + PropKindSize = 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 aecef9347391..1cc1b39f5c4d 100644 --- a/sycl/include/sycl/ext/oneapi/virtual_mem/physical_mem.hpp +++ b/sycl/include/sycl/ext/oneapi/virtual_mem/physical_mem.hpp @@ -13,6 +13,7 @@ #include #include #include +#include #include namespace sycl { @@ -26,17 +27,42 @@ namespace ext::oneapi::experimental { enum class address_access_mode : char { none = 0, read = 1, read_write = 2 }; +// Property controlling whether a physical_mem object can be shared across +// processes via IPC. +struct enable_ipc + : detail::run_time_property_key { + enable_ipc(bool enable = true) : value(enable) {} + bool value; +}; + +using enable_ipc_key = enable_ipc; + +inline bool operator==(const enable_ipc &lhs, const enable_ipc &rhs) { + return lhs.value == rhs.value; +} +inline bool operator!=(const enable_ipc &lhs, const enable_ipc &rhs) { + return !(lhs == rhs); +} + 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); - - physical_mem(const queue &SyclQueue, size_t NumBytes) - : physical_mem(SyclQueue.get_device(), SyclQueue.get_context(), - NumBytes) {} + size_t NumBytes, PropertyListT Props = {}) + : physical_mem(SyclDevice, SyclContext, NumBytes, [&Props]() -> bool { + if constexpr (PropertyListT::template has_property()) + return Props.template get_property().value; + return false; + }()) {} + + template + physical_mem(const queue &SyclQueue, size_t NumBytes, + PropertyListT Props = {}) + : physical_mem(SyclQueue.get_device(), SyclQueue.get_context(), NumBytes, + Props) {} physical_mem(const physical_mem &rhs) = default; physical_mem(physical_mem &&rhs) = default; @@ -57,10 +83,26 @@ class __SYCL_EXPORT physical_mem size_t size() const noexcept; + bool ext_oneapi_ipc_enabled() const; + bool ipc_enabled() const { return ext_oneapi_ipc_enabled(); } + private: + // Internal constructor called by the public templated constructors. + __SYCL_EXPORT physical_mem(const device &SyclDevice, + const context &SyclContext, size_t NumBytes, + bool EnableIpc); + + // Internal constructor for creating a physical_mem from an existing impl + // (used by createSyclObjFromImpl, e.g. when opening from an IPC handle). + explicit physical_mem(std::shared_ptr Impl) + : impl(std::move(Impl)) {} + std::shared_ptr impl; }; +template <> +struct is_property_key_of : std::true_type {}; + } // namespace ext::oneapi::experimental } // namespace _V1 } // namespace sycl diff --git a/sycl/include/sycl/info/aspects.def b/sycl/include/sycl/info/aspects.def index 78785ccf8d64..69ff4e02ee76 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/detail/physical_mem_impl.hpp b/sycl/source/detail/physical_mem_impl.hpp index 3bc4d6865187..3fb19d88b557 100644 --- a/sycl/source/detail/physical_mem_impl.hpp +++ b/sycl/source/detail/physical_mem_impl.hpp @@ -38,9 +38,9 @@ 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), MEnabledIpc(EnableIpc) { adapter_impl &Adapter = MContext->getAdapter(); auto Err = Adapter.call_nocheck( @@ -55,9 +55,22 @@ class physical_mem_impl { Adapter.checkUrResult(Err); } + // Constructor used when opening a physical memory object from an IPC handle. + 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), + MEnabledIpc(false), MOpenedFromIpc(true) {} + ~physical_mem_impl() noexcept(false) { adapter_impl &Adapter = MContext->getAdapter(); - Adapter.call(MPhysicalMem); + // Handles opened via urIPCOpenPhysMemHandleExp must be released with + // urIPCClosePhysMemHandleExp, not urPhysicalMemRelease. + if (MOpenedFromIpc) + Adapter.call( + MContext->getHandleRef(), MPhysicalMem); + else + Adapter.call(MPhysicalMem); } void *map(uintptr_t Ptr, size_t NumBytes, @@ -77,6 +90,7 @@ class physical_mem_impl { } device get_device() const { return createSyclObjFromImpl(MDevice); } size_t size() const noexcept { return MNumBytes; } + bool isEnabledIpc() const noexcept { return MEnabledIpc; } ur_physical_mem_handle_t &getHandleRef() { return MPhysicalMem; } const ur_physical_mem_handle_t &getHandleRef() const { return MPhysicalMem; } @@ -86,6 +100,8 @@ class physical_mem_impl { device_impl &MDevice; const std::shared_ptr MContext; const size_t MNumBytes; + bool MEnabledIpc; + bool MOpenedFromIpc = false; }; } // namespace detail diff --git a/sycl/source/ipc_memory.cpp b/sycl/source/ipc_memory.cpp index 12b47b59d134..89af391164af 100644 --- a/sycl/source/ipc_memory.cpp +++ b/sycl/source/ipc_memory.cpp @@ -8,6 +8,7 @@ #include #include +#include #include #include #include @@ -57,6 +58,42 @@ __SYCL_EXPORT void *openIPCMemHandle(const std::byte *HandleData, return Ptr; } +__SYCL_EXPORT ext::oneapi::experimental::physical_mem +openIPCPhysMemHandle(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(); + + std::byte *NonConstHandleData = const_cast(HandleData); + + ur_physical_mem_handle_t PhysMemHandle = nullptr; + ur_result_t UrRes = + Adapter.call_nocheck( + CtxImpl->getHandleRef(), getSyclObjImpl(Dev)->getHandleRef(), + NonConstHandleData, 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."); + + auto PhysMemImpl = std::make_shared( + *getSyclObjImpl(Dev), Ctx, /*NumBytes=*/0, PhysMemHandle); + return sycl::detail::createSyclObjFromImpl< + ext::oneapi::experimental::physical_mem>(PhysMemImpl); +} + } // namespace detail namespace ext::oneapi::experimental::ipc::memory { @@ -85,6 +122,11 @@ std::pair get(void *Ptr, const sycl::context &Ctx) { CheckDeviceSupport(); Adapter.checkUrResult(UrRes); } + if (HandlePtr == nullptr) + throw sycl::exception( + sycl::make_error_code(errc::runtime), + "urIPCGetMemHandleExp returned success but did not produce a " + "valid IPC handle."); return {HandlePtr, HandleSize}; } @@ -116,6 +158,80 @@ __SYCL_EXPORT void close(void *Ptr, const sycl::context &Ctx) { } } // namespace ext::oneapi::experimental::ipc::memory +namespace ext::oneapi::experimental::ipc::physical_memory { +namespace detail { + +std::pair +get(const ext::oneapi::experimental::physical_mem &PhysMem) { + if (!PhysMem.ext_oneapi_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 CheckDeviceSupport = [&PhysMem]() { + sycl::device Dev = PhysMem.get_device(); + 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 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 = + Adapter.call_nocheck( + CtxImpl->getHandleRef(), PhysMemImpl->getHandleRef(), &HandlePtr, + &HandleSize); + if (UrRes != UR_RESULT_SUCCESS) { + CheckDeviceSupport(); + Adapter.checkUrResult(UrRes); + } + if (HandlePtr == nullptr) + throw sycl::exception( + sycl::make_error_code(errc::runtime), + "urIPCGetPhysMemHandleExp returned success but did not produce a " + "valid IPC handle."); + return {HandlePtr, HandleSize}; +} + +void put(std::byte *HandleData, const sycl::context &Ctx) { + auto CtxImpl = sycl::detail::getSyclObjImpl(Ctx); + CtxImpl->getAdapter().call( + CtxImpl->getHandleRef(), HandleData); +} + +ext::oneapi::experimental::physical_mem open(const std::byte *HandleData, + size_t HandleDataSize, + const sycl::context &Ctx, + const sycl::device &Dev) { + return sycl::detail::openIPCPhysMemHandle(HandleData, HandleDataSize, Ctx, + Dev); +} + +} // namespace detail + +__SYCL_EXPORT ipc::handle +get(const ext::oneapi::experimental::physical_mem &PhysMem) { + std::pair RetHandle = detail::get(PhysMem); + return {RetHandle.first, RetHandle.second}; +} + +__SYCL_EXPORT void put(ipc::handle &Handle, const sycl::context &Ctx) { + detail::put(Handle.MData, Ctx); +} + +__SYCL_EXPORT ext::oneapi::experimental::physical_mem +open(const ipc::handle_data_t &HandleData, const sycl::context &Ctx, + const sycl::device &Dev) { + return detail::open(HandleData.data(), HandleData.size(), Ctx, Dev); +} +} // namespace ext::oneapi::experimental::ipc::physical_memory + namespace ext::oneapi::experimental::ipc_memory { __SYCL_SUPPRESS_DEPRECATED_PUSH __SYCL_EXPORT handle get(void *Ptr, const sycl::context &Ctx) { diff --git a/sycl/source/physical_mem.cpp b/sycl/source/physical_mem.cpp index 67486c83df31..c60ba3eed09c 100644 --- a/sycl/source/physical_mem.cpp +++ b/sycl/source/physical_mem.cpp @@ -14,14 +14,19 @@ inline namespace _V1 { namespace ext::oneapi::experimental { physical_mem::physical_mem(const device &SyclDevice, const context &SyclContext, - size_t NumBytes) { + 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); + *detail::getSyclObjImpl(SyclDevice), SyclContext, NumBytes, EnableIpc); } void *physical_mem::map(uintptr_t Ptr, size_t NumBytes, @@ -32,6 +37,9 @@ 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::ext_oneapi_ipc_enabled() const { + return impl->isEnabledIpc(); +} } // namespace ext::oneapi::experimental } // namespace _V1 diff --git a/unified-runtime/source/adapters/cuda/memory.cpp b/unified-runtime/source/adapters/cuda/memory.cpp index 81b86d744538..0c31094d2f21 100644 --- a/unified-runtime/source/adapters/cuda/memory.cpp +++ b/unified-runtime/source/adapters/cuda/memory.cpp @@ -639,3 +639,24 @@ UR_APIEXPORT ur_result_t UR_APICALL urIPCCloseMemHandleExp(ur_context_handle_t, void *pMem) { return umf::umf2urResult(umfCloseIPCHandle(pMem)); } + +UR_APIEXPORT ur_result_t UR_APICALL urIPCGetPhysMemHandleExp( + ur_context_handle_t, ur_physical_mem_handle_t, void **, size_t *) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urIPCPutPhysMemHandleExp(ur_context_handle_t, void *) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urIPCOpenPhysMemHandleExp(ur_context_handle_t, ur_device_handle_t, void *, + size_t, ur_physical_mem_handle_t *) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urIPCClosePhysMemHandleExp(ur_context_handle_t, ur_physical_mem_handle_t) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} diff --git a/unified-runtime/source/adapters/cuda/ur_interface_loader.cpp b/unified-runtime/source/adapters/cuda/ur_interface_loader.cpp index cfff918f9030..61eac46cd1e7 100644 --- a/unified-runtime/source/adapters/cuda/ur_interface_loader.cpp +++ b/unified-runtime/source/adapters/cuda/ur_interface_loader.cpp @@ -509,6 +509,10 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetIPCExpProcAddrTable( pDdiTable->pfnPutMemHandleExp = urIPCPutMemHandleExp; pDdiTable->pfnOpenMemHandleExp = urIPCOpenMemHandleExp; pDdiTable->pfnCloseMemHandleExp = urIPCCloseMemHandleExp; + pDdiTable->pfnGetPhysMemHandleExp = urIPCGetPhysMemHandleExp; + pDdiTable->pfnPutPhysMemHandleExp = urIPCPutPhysMemHandleExp; + pDdiTable->pfnOpenPhysMemHandleExp = urIPCOpenPhysMemHandleExp; + pDdiTable->pfnClosePhysMemHandleExp = urIPCClosePhysMemHandleExp; return UR_RESULT_SUCCESS; } diff --git a/unified-runtime/source/adapters/hip/memory.cpp b/unified-runtime/source/adapters/hip/memory.cpp index 6360ab0f9235..fec81286ab59 100644 --- a/unified-runtime/source/adapters/hip/memory.cpp +++ b/unified-runtime/source/adapters/hip/memory.cpp @@ -658,3 +658,24 @@ UR_APIEXPORT ur_result_t UR_APICALL urIPCCloseMemHandleExp(ur_context_handle_t, void *) { return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } + +UR_APIEXPORT ur_result_t UR_APICALL urIPCGetPhysMemHandleExp( + ur_context_handle_t, ur_physical_mem_handle_t, void **, size_t *) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urIPCPutPhysMemHandleExp(ur_context_handle_t, void *) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urIPCOpenPhysMemHandleExp(ur_context_handle_t, ur_device_handle_t, void *, + size_t, ur_physical_mem_handle_t *) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urIPCClosePhysMemHandleExp(ur_context_handle_t, ur_physical_mem_handle_t) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} diff --git a/unified-runtime/source/adapters/hip/ur_interface_loader.cpp b/unified-runtime/source/adapters/hip/ur_interface_loader.cpp index 8152546927bf..268aa349b1fc 100644 --- a/unified-runtime/source/adapters/hip/ur_interface_loader.cpp +++ b/unified-runtime/source/adapters/hip/ur_interface_loader.cpp @@ -502,6 +502,10 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetIPCExpProcAddrTable( pDdiTable->pfnPutMemHandleExp = urIPCPutMemHandleExp; pDdiTable->pfnOpenMemHandleExp = urIPCOpenMemHandleExp; pDdiTable->pfnCloseMemHandleExp = urIPCCloseMemHandleExp; + pDdiTable->pfnGetPhysMemHandleExp = urIPCGetPhysMemHandleExp; + pDdiTable->pfnPutPhysMemHandleExp = urIPCPutPhysMemHandleExp; + pDdiTable->pfnOpenPhysMemHandleExp = urIPCOpenPhysMemHandleExp; + pDdiTable->pfnClosePhysMemHandleExp = urIPCClosePhysMemHandleExp; return UR_RESULT_SUCCESS; } diff --git a/unified-runtime/source/adapters/level_zero/memory.cpp b/unified-runtime/source/adapters/level_zero/memory.cpp index f0bdf7d23e59..4a02d167a53f 100644 --- a/unified-runtime/source/adapters/level_zero/memory.cpp +++ b/unified-runtime/source/adapters/level_zero/memory.cpp @@ -2011,6 +2011,27 @@ ur_result_t urIPCCloseMemHandleExp(ur_context_handle_t, void *pMem) { return umf::umf2urResult(umfCloseIPCHandle(pMem)); } +ur_result_t urIPCGetPhysMemHandleExp(ur_context_handle_t, + ur_physical_mem_handle_t, void **, + size_t *) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +ur_result_t urIPCPutPhysMemHandleExp(ur_context_handle_t, void *) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +ur_result_t urIPCOpenPhysMemHandleExp(ur_context_handle_t, ur_device_handle_t, + void *, size_t, + ur_physical_mem_handle_t *) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +ur_result_t urIPCClosePhysMemHandleExp(ur_context_handle_t, + ur_physical_mem_handle_t) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + } // namespace ur::level_zero // If indirect access tracking is enabled then performs reference counting, diff --git a/unified-runtime/source/adapters/level_zero/v2/memory.cpp b/unified-runtime/source/adapters/level_zero/v2/memory.cpp index b0601ba956af..c006df99ed11 100644 --- a/unified-runtime/source/adapters/level_zero/v2/memory.cpp +++ b/unified-runtime/source/adapters/level_zero/v2/memory.cpp @@ -925,4 +925,25 @@ ur_result_t urIPCCloseMemHandleExp(ur_context_handle_t, void *pMem) { return umf::umf2urResult(umfCloseIPCHandle(pMem)); } +ur_result_t urIPCGetPhysMemHandleExp(ur_context_handle_t, + ur_physical_mem_handle_t, void **, + size_t *) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +ur_result_t urIPCPutPhysMemHandleExp(ur_context_handle_t, void *) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +ur_result_t urIPCOpenPhysMemHandleExp(ur_context_handle_t, ur_device_handle_t, + void *, size_t, + ur_physical_mem_handle_t *) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +ur_result_t urIPCClosePhysMemHandleExp(ur_context_handle_t, + ur_physical_mem_handle_t) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + } // namespace ur::level_zero diff --git a/unified-runtime/source/adapters/native_cpu/memory.cpp b/unified-runtime/source/adapters/native_cpu/memory.cpp index 2ac5eccc11cc..0d77178b8890 100644 --- a/unified-runtime/source/adapters/native_cpu/memory.cpp +++ b/unified-runtime/source/adapters/native_cpu/memory.cpp @@ -160,3 +160,24 @@ UR_APIEXPORT ur_result_t UR_APICALL urIPCCloseMemHandleExp(ur_context_handle_t, void *) { return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } + +UR_APIEXPORT ur_result_t UR_APICALL urIPCGetPhysMemHandleExp( + ur_context_handle_t, ur_physical_mem_handle_t, void **, size_t *) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urIPCPutPhysMemHandleExp(ur_context_handle_t, void *) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urIPCOpenPhysMemHandleExp(ur_context_handle_t, ur_device_handle_t, void *, + size_t, ur_physical_mem_handle_t *) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urIPCClosePhysMemHandleExp(ur_context_handle_t, ur_physical_mem_handle_t) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} diff --git a/unified-runtime/source/adapters/native_cpu/ur_interface_loader.cpp b/unified-runtime/source/adapters/native_cpu/ur_interface_loader.cpp index bc1183e1c9ce..ce2bb819bec7 100644 --- a/unified-runtime/source/adapters/native_cpu/ur_interface_loader.cpp +++ b/unified-runtime/source/adapters/native_cpu/ur_interface_loader.cpp @@ -486,6 +486,10 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetIPCExpProcAddrTable( pDdiTable->pfnPutMemHandleExp = urIPCPutMemHandleExp; pDdiTable->pfnOpenMemHandleExp = urIPCOpenMemHandleExp; pDdiTable->pfnCloseMemHandleExp = urIPCCloseMemHandleExp; + pDdiTable->pfnGetPhysMemHandleExp = urIPCGetPhysMemHandleExp; + pDdiTable->pfnPutPhysMemHandleExp = urIPCPutPhysMemHandleExp; + pDdiTable->pfnOpenPhysMemHandleExp = urIPCOpenPhysMemHandleExp; + pDdiTable->pfnClosePhysMemHandleExp = urIPCClosePhysMemHandleExp; return UR_RESULT_SUCCESS; } diff --git a/unified-runtime/source/adapters/offload/memory.cpp b/unified-runtime/source/adapters/offload/memory.cpp index c0b7a97fb62e..183f2e82e9ba 100644 --- a/unified-runtime/source/adapters/offload/memory.cpp +++ b/unified-runtime/source/adapters/offload/memory.cpp @@ -229,3 +229,24 @@ UR_APIEXPORT ur_result_t UR_APICALL urIPCCloseMemHandleExp(ur_context_handle_t, void *) { return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } + +UR_APIEXPORT ur_result_t UR_APICALL urIPCGetPhysMemHandleExp( + ur_context_handle_t, ur_physical_mem_handle_t, void **, size_t *) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urIPCPutPhysMemHandleExp(ur_context_handle_t, void *) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urIPCOpenPhysMemHandleExp(ur_context_handle_t, ur_device_handle_t, void *, + size_t, ur_physical_mem_handle_t *) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urIPCClosePhysMemHandleExp(ur_context_handle_t, ur_physical_mem_handle_t) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} diff --git a/unified-runtime/source/adapters/offload/ur_interface_loader.cpp b/unified-runtime/source/adapters/offload/ur_interface_loader.cpp index 69d9aafaf4a6..ff6fbc864547 100644 --- a/unified-runtime/source/adapters/offload/ur_interface_loader.cpp +++ b/unified-runtime/source/adapters/offload/ur_interface_loader.cpp @@ -436,6 +436,10 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetIPCExpProcAddrTable( pDdiTable->pfnPutMemHandleExp = urIPCPutMemHandleExp; pDdiTable->pfnOpenMemHandleExp = urIPCOpenMemHandleExp; pDdiTable->pfnCloseMemHandleExp = urIPCCloseMemHandleExp; + pDdiTable->pfnGetPhysMemHandleExp = urIPCGetPhysMemHandleExp; + pDdiTable->pfnPutPhysMemHandleExp = urIPCPutPhysMemHandleExp; + pDdiTable->pfnOpenPhysMemHandleExp = urIPCOpenPhysMemHandleExp; + pDdiTable->pfnClosePhysMemHandleExp = urIPCClosePhysMemHandleExp; return UR_RESULT_SUCCESS; } diff --git a/unified-runtime/source/adapters/opencl/memory.cpp b/unified-runtime/source/adapters/opencl/memory.cpp index ca9e88d4f1ab..dbcfe573ecac 100644 --- a/unified-runtime/source/adapters/opencl/memory.cpp +++ b/unified-runtime/source/adapters/opencl/memory.cpp @@ -601,3 +601,24 @@ UR_APIEXPORT ur_result_t UR_APICALL urIPCCloseMemHandleExp(ur_context_handle_t, void *) { return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } + +UR_APIEXPORT ur_result_t UR_APICALL urIPCGetPhysMemHandleExp( + ur_context_handle_t, ur_physical_mem_handle_t, void **, size_t *) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urIPCPutPhysMemHandleExp(ur_context_handle_t, void *) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urIPCOpenPhysMemHandleExp(ur_context_handle_t, ur_device_handle_t, void *, + size_t, ur_physical_mem_handle_t *) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urIPCClosePhysMemHandleExp(ur_context_handle_t, ur_physical_mem_handle_t) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} diff --git a/unified-runtime/source/adapters/opencl/ur_interface_loader.cpp b/unified-runtime/source/adapters/opencl/ur_interface_loader.cpp index 8e1c9ee94b8c..a7dbb9a152a1 100644 --- a/unified-runtime/source/adapters/opencl/ur_interface_loader.cpp +++ b/unified-runtime/source/adapters/opencl/ur_interface_loader.cpp @@ -458,6 +458,10 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetIPCExpProcAddrTable( pDdiTable->pfnPutMemHandleExp = urIPCPutMemHandleExp; pDdiTable->pfnOpenMemHandleExp = urIPCOpenMemHandleExp; pDdiTable->pfnCloseMemHandleExp = urIPCCloseMemHandleExp; + pDdiTable->pfnGetPhysMemHandleExp = urIPCGetPhysMemHandleExp; + pDdiTable->pfnPutPhysMemHandleExp = urIPCPutPhysMemHandleExp; + pDdiTable->pfnOpenPhysMemHandleExp = urIPCOpenPhysMemHandleExp; + pDdiTable->pfnClosePhysMemHandleExp = urIPCClosePhysMemHandleExp; return UR_RESULT_SUCCESS; } From 5ebb852cf446cd74225329227c7cc1713919c8a3 Mon Sep 17 00:00:00 2001 From: Lukasz Dorau Date: Thu, 21 May 2026 09:57:57 +0000 Subject: [PATCH 05/17] [UR] Add UR_PHYSICAL_MEM_FLAG_ENABLE_IPC to ur_physical_mem_flags_t Add UR_PHYSICAL_MEM_FLAG_ENABLE_IPC = UR_BIT(1) to ur_physical_mem_flags_t in virtual_memory.yml. This flag allows callers to request IPC-exportable physical memory allocations at creation time via urPhysicalMemCreate. Signed-off-by: Lukasz Dorau --- unified-runtime/scripts/core/virtual_memory.yml | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/unified-runtime/scripts/core/virtual_memory.yml b/unified-runtime/scripts/core/virtual_memory.yml index 00d03c54a572..4fe51b74dd55 100644 --- a/unified-runtime/scripts/core/virtual_memory.yml +++ b/unified-runtime/scripts/core/virtual_memory.yml @@ -246,8 +246,9 @@ desc: "Physical memory creation properties." class: $xPhysicalMem name: $x_physical_mem_flags_t etors: - - name: TBD - desc: "reserved for future use." + - name: ENABLE_IPC + value: "$X_BIT(1)" + desc: "allocate physical memory that can be shared via IPC handles." --- #-------------------------------------------------------------------------- type: struct From 936049e9486935f2c7e81ba31513ad5dbd75d365 Mon Sep 17 00:00:00 2001 From: Lukasz Dorau Date: Thu, 21 May 2026 10:24:31 +0000 Subject: [PATCH 06/17] [UR] Regenerate headers for UR_PHYSICAL_MEM_FLAG_ENABLE_IPC Regenerate ur_api.h and ur_print.hpp to include the new UR_PHYSICAL_MEM_FLAG_ENABLE_IPC flag and update UR_PHYSICAL_MEM_FLAGS_MASK from 0xfffffffe to 0xfffffffc. Signed-off-by: Lukasz Dorau --- unified-runtime/include/unified-runtime/ur_api.h | 6 +++--- unified-runtime/include/unified-runtime/ur_print.hpp | 11 ++++++----- 2 files changed, 9 insertions(+), 8 deletions(-) diff --git a/unified-runtime/include/unified-runtime/ur_api.h b/unified-runtime/include/unified-runtime/ur_api.h index 826feace5c79..6c252e35a4c9 100644 --- a/unified-runtime/include/unified-runtime/ur_api.h +++ b/unified-runtime/include/unified-runtime/ur_api.h @@ -5342,15 +5342,15 @@ UR_APIEXPORT ur_result_t UR_APICALL urVirtualMemGetInfo( /// @brief Physical memory creation properties. typedef uint32_t ur_physical_mem_flags_t; typedef enum ur_physical_mem_flag_t { - /// reserved for future use. - UR_PHYSICAL_MEM_FLAG_TBD = UR_BIT(0), + /// allocate physical memory that can be shared via IPC handles. + UR_PHYSICAL_MEM_FLAG_ENABLE_IPC = UR_BIT(1), /// @cond UR_PHYSICAL_MEM_FLAG_FORCE_UINT32 = 0x7fffffff /// @endcond } ur_physical_mem_flag_t; /// @brief Bit Mask for validating ur_physical_mem_flags_t -#define UR_PHYSICAL_MEM_FLAGS_MASK 0xfffffffe +#define UR_PHYSICAL_MEM_FLAGS_MASK 0xfffffffd /////////////////////////////////////////////////////////////////////////////// /// @brief Physical memory creation properties. diff --git a/unified-runtime/include/unified-runtime/ur_print.hpp b/unified-runtime/include/unified-runtime/ur_print.hpp index d0728babf7a6..50a876b67ca6 100644 --- a/unified-runtime/include/unified-runtime/ur_print.hpp +++ b/unified-runtime/include/unified-runtime/ur_print.hpp @@ -9153,8 +9153,8 @@ inline ur_result_t printTagged(std::ostream &os, const void *ptr, inline std::ostream &operator<<(std::ostream &os, enum ur_physical_mem_flag_t value) { switch (value) { - case UR_PHYSICAL_MEM_FLAG_TBD: - os << "UR_PHYSICAL_MEM_FLAG_TBD"; + case UR_PHYSICAL_MEM_FLAG_ENABLE_IPC: + os << "UR_PHYSICAL_MEM_FLAG_ENABLE_IPC"; break; default: os << "unknown enumerator"; @@ -9172,14 +9172,15 @@ inline ur_result_t printFlag(std::ostream &os, uint32_t val = flag; bool first = true; - if ((val & UR_PHYSICAL_MEM_FLAG_TBD) == (uint32_t)UR_PHYSICAL_MEM_FLAG_TBD) { - val ^= (uint32_t)UR_PHYSICAL_MEM_FLAG_TBD; + if ((val & UR_PHYSICAL_MEM_FLAG_ENABLE_IPC) == + (uint32_t)UR_PHYSICAL_MEM_FLAG_ENABLE_IPC) { + val ^= (uint32_t)UR_PHYSICAL_MEM_FLAG_ENABLE_IPC; if (!first) { os << " | "; } else { first = false; } - os << UR_PHYSICAL_MEM_FLAG_TBD; + os << UR_PHYSICAL_MEM_FLAG_ENABLE_IPC; } if (val != 0) { std::bitset<32> bits(val); From c3acf0817a4be20b61932f550ae16f855a410a25 Mon Sep 17 00:00:00 2001 From: Lukasz Dorau Date: Thu, 21 May 2026 09:01:57 +0000 Subject: [PATCH 07/17] [UR][L0] Implement IPC API for physical_mem objects Implement the four experimental IPC functions for physical memory objects in the Level Zero (v1 and v2) adapters: - urIPCGetPhysMemHandleExp - urIPCPutPhysMemHandleExp - urIPCOpenPhysMemHandleExp - urIPCClosePhysMemHandleExp Level Zero adapter changes (shared by v1 and v2): - physical_mem.hpp: add Size member to ur_physical_mem_handle_t_ (needed on the import side); add ZeIPCPhysMemHandleData struct that holds a ze_ipc_mem_handle_t and the allocation size (72 bytes total) - physical_mem.cpp: chain ze_external_memory_export_desc_t into ze_physical_mem_desc_t.pNext when UR_PHYSICAL_MEM_FLAG_ENABLE_IPC is set, so the allocation is IPC-exportable; fix ZePhysicalMem leak in exception handlers of urPhysicalMemCreate - memory.cpp and v2/memory.cpp: replace UNSUPPORTED_FEATURE stubs with full implementations using zePhysicalMemGetProperties, zeMemGetIpcHandleFromFileDescriptorExp, zeMemGetFileDescriptorFromIpcHandleExp, zeMemPutIpcHandle, zePhysicalMemCreate, zePhysicalMemDestroy SYCL layer changes: - physical_mem_impl.hpp: pass UR_PHYSICAL_MEM_FLAG_ENABLE_IPC in ur_physical_mem_properties_t when EnableIpc is true, so the adapter creates an exportable allocation Signed-off-by: Lukasz Dorau --- sycl/source/detail/physical_mem_impl.hpp | 7 +- .../source/adapters/level_zero/memory.cpp | 122 +++++++++++++++-- .../adapters/level_zero/physical_mem.cpp | 16 ++- .../adapters/level_zero/physical_mem.hpp | 15 ++- .../source/adapters/level_zero/v2/memory.cpp | 123 ++++++++++++++++-- 5 files changed, 253 insertions(+), 30 deletions(-) diff --git a/sycl/source/detail/physical_mem_impl.hpp b/sycl/source/detail/physical_mem_impl.hpp index 3fb19d88b557..e330a13e8faa 100644 --- a/sycl/source/detail/physical_mem_impl.hpp +++ b/sycl/source/detail/physical_mem_impl.hpp @@ -43,8 +43,13 @@ class physical_mem_impl { MNumBytes(NumBytes), MEnabledIpc(EnableIpc) { adapter_impl &Adapter = MContext->getAdapter(); + ur_physical_mem_properties_t Props = { + UR_STRUCTURE_TYPE_PHYSICAL_MEM_PROPERTIES, nullptr, + EnableIpc ? UR_PHYSICAL_MEM_FLAG_ENABLE_IPC + : ur_physical_mem_flags_t(0)}; + auto Err = Adapter.call_nocheck( - MContext->getHandleRef(), MDevice.getHandleRef(), MNumBytes, nullptr, + MContext->getHandleRef(), MDevice.getHandleRef(), MNumBytes, &Props, &MPhysicalMem); if (Err == UR_RESULT_ERROR_OUT_OF_RESOURCES || diff --git a/unified-runtime/source/adapters/level_zero/memory.cpp b/unified-runtime/source/adapters/level_zero/memory.cpp index 4a02d167a53f..29c54881e58b 100644 --- a/unified-runtime/source/adapters/level_zero/memory.cpp +++ b/unified-runtime/source/adapters/level_zero/memory.cpp @@ -10,6 +10,7 @@ #include #include #include +#include #include #include "context.hpp" @@ -17,6 +18,7 @@ #include "helpers/memory_helpers.hpp" #include "image_common.hpp" #include "logger/ur_logger.hpp" +#include "physical_mem.hpp" #include "queue.hpp" #include "ur_interface_loader.hpp" #include "ur_level_zero.hpp" @@ -2011,25 +2013,119 @@ ur_result_t urIPCCloseMemHandleExp(ur_context_handle_t, void *pMem) { return umf::umf2urResult(umfCloseIPCHandle(pMem)); } -ur_result_t urIPCGetPhysMemHandleExp(ur_context_handle_t, - ur_physical_mem_handle_t, void **, - size_t *) { - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +ur_result_t urIPCGetPhysMemHandleExp(ur_context_handle_t hContext, + ur_physical_mem_handle_t hPhysMem, + void **ppIPCPhysMemHandleData, + size_t *pIPCPhysMemHandleDataSizeRet) { + // Fast path: size-only query. + if (!ppIPCPhysMemHandleData) { + if (pIPCPhysMemHandleDataSizeRet) + *pIPCPhysMemHandleDataSizeRet = sizeof(ZeIPCPhysMemHandleData); + return UR_RESULT_SUCCESS; + } + + // Export the physical memory object as an opaque file descriptor. + ze_external_memory_export_fd_t ExportFd = {}; + ExportFd.stype = ZE_STRUCTURE_TYPE_EXTERNAL_MEMORY_EXPORT_FD; + ExportFd.pNext = nullptr; + ExportFd.flags = ZE_EXTERNAL_MEMORY_TYPE_FLAG_OPAQUE_FD; + + ze_physical_mem_properties_t Props = {}; + Props.stype = ZE_STRUCTURE_TYPE_PHYSICAL_MEM_PROPERTIES; + Props.pNext = &ExportFd; + + ZE2UR_CALL(zePhysicalMemGetProperties, + (hContext->getZeHandle(), hPhysMem->ZePhysicalMem, &Props)); + + // Convert the file descriptor to an opaque Level Zero IPC handle so that + // the caller can pass it across process boundaries. + ze_ipc_mem_handle_t ZeIpcHandle = {}; + ze_result_t ZeRes = zeMemGetIpcHandleFromFileDescriptorExp( + hContext->getZeHandle(), static_cast(ExportFd.fd), + &ZeIpcHandle); + // The exported fd is no longer needed once the IPC handle is obtained. + close(ExportFd.fd); + if (ZeRes != ZE_RESULT_SUCCESS) + return ze2urResult(ZeRes); + + auto *HandleData = new (std::nothrow) ZeIPCPhysMemHandleData; + if (!HandleData) { + zeMemPutIpcHandle(hContext->getZeHandle(), ZeIpcHandle); + return UR_RESULT_ERROR_OUT_OF_HOST_MEMORY; + } + + HandleData->ZeHandle = ZeIpcHandle; + HandleData->Size = hPhysMem->Size; + + *ppIPCPhysMemHandleData = HandleData; + if (pIPCPhysMemHandleDataSizeRet) + *pIPCPhysMemHandleDataSizeRet = sizeof(ZeIPCPhysMemHandleData); + return UR_RESULT_SUCCESS; } -ur_result_t urIPCPutPhysMemHandleExp(ur_context_handle_t, void *) { - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +ur_result_t urIPCPutPhysMemHandleExp(ur_context_handle_t hContext, + void *pIPCPhysMemHandleData) { + auto *HandleData = + static_cast(pIPCPhysMemHandleData); + ur_result_t Res = ze2urResult( + zeMemPutIpcHandle(hContext->getZeHandle(), HandleData->ZeHandle)); + delete HandleData; + return Res; } -ur_result_t urIPCOpenPhysMemHandleExp(ur_context_handle_t, ur_device_handle_t, - void *, size_t, - ur_physical_mem_handle_t *) { - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +ur_result_t urIPCOpenPhysMemHandleExp(ur_context_handle_t hContext, + ur_device_handle_t hDevice, + void *pIPCPhysMemHandleData, + size_t ipcPhysMemHandleDataSize, + ur_physical_mem_handle_t *phPhysMem) { + if (ipcPhysMemHandleDataSize != sizeof(ZeIPCPhysMemHandleData)) + return UR_RESULT_ERROR_INVALID_VALUE; + + auto *HandleData = + static_cast(pIPCPhysMemHandleData); + + // Convert the opaque IPC handle back to a file descriptor for import. + uint64_t Fd = 0; + ZE2UR_CALL(zeMemGetFileDescriptorFromIpcHandleExp, + (hContext->getZeHandle(), HandleData->ZeHandle, &Fd)); + + ze_external_memory_import_fd_t ImportFd = {}; + ImportFd.stype = ZE_STRUCTURE_TYPE_EXTERNAL_MEMORY_IMPORT_FD; + ImportFd.pNext = nullptr; + ImportFd.flags = ZE_EXTERNAL_MEMORY_TYPE_FLAG_OPAQUE_FD; + ImportFd.fd = static_cast(Fd); + + ZeStruct PhysMemDesc; + PhysMemDesc.pNext = &ImportFd; + PhysMemDesc.flags = 0; + PhysMemDesc.size = HandleData->Size; + + ze_physical_mem_handle_t ZePhysMem; + ze_result_t ZeRes = zePhysicalMemCreate( + hContext->getZeHandle(), hDevice->ZeDevice, &PhysMemDesc, &ZePhysMem); + // The imported fd is no longer needed once zePhysicalMemCreate returns. + close(static_cast(Fd)); + if (ZeRes != ZE_RESULT_SUCCESS) + return ze2urResult(ZeRes); + try { + *phPhysMem = + new ur_physical_mem_handle_t_(ZePhysMem, hContext, HandleData->Size); + } catch (const std::bad_alloc &) { + zePhysicalMemDestroy(hContext->getZeHandle(), ZePhysMem); + return UR_RESULT_ERROR_OUT_OF_HOST_MEMORY; + } catch (...) { + zePhysicalMemDestroy(hContext->getZeHandle(), ZePhysMem); + return UR_RESULT_ERROR_UNKNOWN; + } + return UR_RESULT_SUCCESS; } -ur_result_t urIPCClosePhysMemHandleExp(ur_context_handle_t, - ur_physical_mem_handle_t) { - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +ur_result_t urIPCClosePhysMemHandleExp(ur_context_handle_t hContext, + ur_physical_mem_handle_t hPhysMem) { + ur_result_t Res = ze2urResult( + zePhysicalMemDestroy(hContext->getZeHandle(), hPhysMem->ZePhysicalMem)); + delete hPhysMem; + return Res; } } // namespace ur::level_zero diff --git a/unified-runtime/source/adapters/level_zero/physical_mem.cpp b/unified-runtime/source/adapters/level_zero/physical_mem.cpp index a0663253f42e..8f3010d1b683 100644 --- a/unified-runtime/source/adapters/level_zero/physical_mem.cpp +++ b/unified-runtime/source/adapters/level_zero/physical_mem.cpp @@ -27,14 +27,28 @@ ur_result_t urPhysicalMemCreate( PhysicalMemDesc.flags = 0; PhysicalMemDesc.size = size; + // If IPC export is requested, chain in the export descriptor so the + // physical memory can later be shared via urIPCGetPhysMemHandleExp. + ze_external_memory_export_desc_t ExportDesc = {}; + ExportDesc.stype = ZE_STRUCTURE_TYPE_EXTERNAL_MEMORY_EXPORT_DESC; + ExportDesc.pNext = nullptr; + ExportDesc.flags = ZE_EXTERNAL_MEMORY_TYPE_FLAG_OPAQUE_FD; + bool EnableIpc = + pProperties && (pProperties->flags & UR_PHYSICAL_MEM_FLAG_ENABLE_IPC); + if (EnableIpc) + PhysicalMemDesc.pNext = &ExportDesc; + ze_physical_mem_handle_t ZePhysicalMem; ZE2UR_CALL(zePhysicalMemCreate, (hContext->getZeHandle(), hDevice->ZeDevice, &PhysicalMemDesc, &ZePhysicalMem)); try { - *phPhysicalMem = new ur_physical_mem_handle_t_(ZePhysicalMem, hContext); + *phPhysicalMem = + new ur_physical_mem_handle_t_(ZePhysicalMem, hContext, size); } catch (const std::bad_alloc &) { + zePhysicalMemDestroy(hContext->getZeHandle(), ZePhysicalMem); return UR_RESULT_ERROR_OUT_OF_HOST_MEMORY; } catch (...) { + zePhysicalMemDestroy(hContext->getZeHandle(), ZePhysicalMem); return UR_RESULT_ERROR_UNKNOWN; } return UR_RESULT_SUCCESS; diff --git a/unified-runtime/source/adapters/level_zero/physical_mem.hpp b/unified-runtime/source/adapters/level_zero/physical_mem.hpp index a5db639575e8..54e5229d09cc 100644 --- a/unified-runtime/source/adapters/level_zero/physical_mem.hpp +++ b/unified-runtime/source/adapters/level_zero/physical_mem.hpp @@ -11,10 +11,18 @@ #include "common.hpp" #include "common/ur_ref_count.hpp" +// Opaque handle data exchanged between processes for physical memory IPC. +// Contains a Level Zero IPC memory handle and the size of the allocation +// (required by zePhysicalMemCreate on the receiving side). +struct ZeIPCPhysMemHandleData { + ze_ipc_mem_handle_t ZeHandle; + size_t Size; +}; + struct ur_physical_mem_handle_t_ : ur_object { ur_physical_mem_handle_t_(ze_physical_mem_handle_t ZePhysicalMem, - ur_context_handle_t Context) - : ZePhysicalMem{ZePhysicalMem}, Context{Context} {} + ur_context_handle_t Context, size_t Size) + : ZePhysicalMem{ZePhysicalMem}, Context{Context}, Size{Size} {} // Level Zero physical memory handle. ze_physical_mem_handle_t ZePhysicalMem; @@ -22,5 +30,8 @@ struct ur_physical_mem_handle_t_ : ur_object { // Keeps the PI context of this memory handle. ur_context_handle_t Context; + // Size in bytes of this physical memory allocation. + size_t Size; + ur::RefCount RefCount; }; diff --git a/unified-runtime/source/adapters/level_zero/v2/memory.cpp b/unified-runtime/source/adapters/level_zero/v2/memory.cpp index c006df99ed11..3541d15a95eb 100644 --- a/unified-runtime/source/adapters/level_zero/v2/memory.cpp +++ b/unified-runtime/source/adapters/level_zero/v2/memory.cpp @@ -9,11 +9,14 @@ #include "memory.hpp" +#include + #include "../ur_interface_loader.hpp" #include "context.hpp" #include "../helpers/memory_helpers.hpp" #include "../image_common.hpp" +#include "../physical_mem.hpp" static bool isAccessCompatible(ur_mem_buffer_t::device_access_mode_t requested, ur_mem_buffer_t::device_access_mode_t actual) { @@ -925,25 +928,119 @@ ur_result_t urIPCCloseMemHandleExp(ur_context_handle_t, void *pMem) { return umf::umf2urResult(umfCloseIPCHandle(pMem)); } -ur_result_t urIPCGetPhysMemHandleExp(ur_context_handle_t, - ur_physical_mem_handle_t, void **, - size_t *) { - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +ur_result_t urIPCGetPhysMemHandleExp(ur_context_handle_t hContext, + ur_physical_mem_handle_t hPhysMem, + void **ppIPCPhysMemHandleData, + size_t *pIPCPhysMemHandleDataSizeRet) { + // Fast path: size-only query. + if (!ppIPCPhysMemHandleData) { + if (pIPCPhysMemHandleDataSizeRet) + *pIPCPhysMemHandleDataSizeRet = sizeof(ZeIPCPhysMemHandleData); + return UR_RESULT_SUCCESS; + } + + // Export the physical memory object as an opaque file descriptor. + ze_external_memory_export_fd_t ExportFd = {}; + ExportFd.stype = ZE_STRUCTURE_TYPE_EXTERNAL_MEMORY_EXPORT_FD; + ExportFd.pNext = nullptr; + ExportFd.flags = ZE_EXTERNAL_MEMORY_TYPE_FLAG_OPAQUE_FD; + + ze_physical_mem_properties_t Props = {}; + Props.stype = ZE_STRUCTURE_TYPE_PHYSICAL_MEM_PROPERTIES; + Props.pNext = &ExportFd; + + ZE2UR_CALL(zePhysicalMemGetProperties, + (hContext->getZeHandle(), hPhysMem->ZePhysicalMem, &Props)); + + // Convert the file descriptor to an opaque Level Zero IPC handle so that + // the caller can pass it across process boundaries. + ze_ipc_mem_handle_t ZeIpcHandle = {}; + ze_result_t ZeRes = zeMemGetIpcHandleFromFileDescriptorExp( + hContext->getZeHandle(), static_cast(ExportFd.fd), + &ZeIpcHandle); + // The exported fd is no longer needed once the IPC handle is obtained. + close(ExportFd.fd); + if (ZeRes != ZE_RESULT_SUCCESS) + return ze2urResult(ZeRes); + + auto *HandleData = new (std::nothrow) ZeIPCPhysMemHandleData; + if (!HandleData) { + zeMemPutIpcHandle(hContext->getZeHandle(), ZeIpcHandle); + return UR_RESULT_ERROR_OUT_OF_HOST_MEMORY; + } + + HandleData->ZeHandle = ZeIpcHandle; + HandleData->Size = hPhysMem->Size; + + *ppIPCPhysMemHandleData = HandleData; + if (pIPCPhysMemHandleDataSizeRet) + *pIPCPhysMemHandleDataSizeRet = sizeof(ZeIPCPhysMemHandleData); + return UR_RESULT_SUCCESS; } -ur_result_t urIPCPutPhysMemHandleExp(ur_context_handle_t, void *) { - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +ur_result_t urIPCPutPhysMemHandleExp(ur_context_handle_t hContext, + void *pIPCPhysMemHandleData) { + auto *HandleData = + static_cast(pIPCPhysMemHandleData); + ur_result_t Res = ze2urResult( + zeMemPutIpcHandle(hContext->getZeHandle(), HandleData->ZeHandle)); + delete HandleData; + return Res; } -ur_result_t urIPCOpenPhysMemHandleExp(ur_context_handle_t, ur_device_handle_t, - void *, size_t, - ur_physical_mem_handle_t *) { - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +ur_result_t urIPCOpenPhysMemHandleExp(ur_context_handle_t hContext, + ur_device_handle_t hDevice, + void *pIPCPhysMemHandleData, + size_t ipcPhysMemHandleDataSize, + ur_physical_mem_handle_t *phPhysMem) { + if (ipcPhysMemHandleDataSize != sizeof(ZeIPCPhysMemHandleData)) + return UR_RESULT_ERROR_INVALID_VALUE; + + auto *HandleData = + static_cast(pIPCPhysMemHandleData); + + // Convert the opaque IPC handle back to a file descriptor for import. + uint64_t Fd = 0; + ZE2UR_CALL(zeMemGetFileDescriptorFromIpcHandleExp, + (hContext->getZeHandle(), HandleData->ZeHandle, &Fd)); + + ze_external_memory_import_fd_t ImportFd = {}; + ImportFd.stype = ZE_STRUCTURE_TYPE_EXTERNAL_MEMORY_IMPORT_FD; + ImportFd.pNext = nullptr; + ImportFd.flags = ZE_EXTERNAL_MEMORY_TYPE_FLAG_OPAQUE_FD; + ImportFd.fd = static_cast(Fd); + + ZeStruct PhysMemDesc; + PhysMemDesc.pNext = &ImportFd; + PhysMemDesc.flags = 0; + PhysMemDesc.size = HandleData->Size; + + ze_physical_mem_handle_t ZePhysMem; + ze_result_t ZeRes = zePhysicalMemCreate( + hContext->getZeHandle(), hDevice->ZeDevice, &PhysMemDesc, &ZePhysMem); + // The imported fd is no longer needed once zePhysicalMemCreate returns. + close(static_cast(Fd)); + if (ZeRes != ZE_RESULT_SUCCESS) + return ze2urResult(ZeRes); + try { + *phPhysMem = + new ur_physical_mem_handle_t_(ZePhysMem, hContext, HandleData->Size); + } catch (const std::bad_alloc &) { + zePhysicalMemDestroy(hContext->getZeHandle(), ZePhysMem); + return UR_RESULT_ERROR_OUT_OF_HOST_MEMORY; + } catch (...) { + zePhysicalMemDestroy(hContext->getZeHandle(), ZePhysMem); + return UR_RESULT_ERROR_UNKNOWN; + } + return UR_RESULT_SUCCESS; } -ur_result_t urIPCClosePhysMemHandleExp(ur_context_handle_t, - ur_physical_mem_handle_t) { - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +ur_result_t urIPCClosePhysMemHandleExp(ur_context_handle_t hContext, + ur_physical_mem_handle_t hPhysMem) { + ur_result_t Res = ze2urResult( + zePhysicalMemDestroy(hContext->getZeHandle(), hPhysMem->ZePhysicalMem)); + delete hPhysMem; + return Res; } } // namespace ur::level_zero From 42fa566f9cad386708782f10f2a0d72b90917afd Mon Sep 17 00:00:00 2001 From: Lukasz Dorau Date: Thu, 21 May 2026 11:09:06 +0000 Subject: [PATCH 08/17] [UR][L0] Add and fix conformance tests for IPC physical_mem UR API Fix urPhysicalMemCreate conformance test: Replace UR_PHYSICAL_MEM_FLAG_TBD (removed from ur_physical_mem_flags_t) with UR_PHYSICAL_MEM_FLAG_ENABLE_IPC in urPhysicalMemCreateWithFlagsParamTest. Add conformance tests for the four experimental IPC physical memory functions: - urIPCGetPhysMemHandleExp - urIPCPutPhysMemHandleExp - urIPCOpenPhysMemHandleExp - urIPCClosePhysMemHandleExp Tests are placed in test/conformance/virtual_memory/ alongside the existing physical memory tests. A shared fixture header urIPCPhysMemHandleExpFixtures.hpp defines three fixture classes: - urIPCPhysMemTest: creates a physical_mem with UR_PHYSICAL_MEM_FLAG_ENABLE_IPC; skips if the device does not support IPC memory (UR_DEVICE_INFO_IPC_MEMORY_SUPPORT_EXP) or if urPhysicalMemCreate returns UR_RESULT_ERROR_UNSUPPORTED_FEATURE - urIPCPhysMemHandleTest: additionally acquires an IPC handle via urIPCGetPhysMemHandleExp; skips on UR_RESULT_ERROR_UNSUPPORTED_FEATURE - urIPCOpenedPhysMemTest: additionally opens the IPC handle via urIPCOpenPhysMemHandleExp Also add a standalone Level Zero e2e diagnostic test ze_ipc_phys_mem_e2e.cpp that exercises the full IPC flow directly through the L0 API (not part of the CMake build; build instructions are in the file header comment). Fix L0 adapter IPC phys_mem handle implementation: Replace the zeMemGetIpcHandleFromFileDescriptorExp / zeMemGetFileDescriptorFromIpcHandleExp round-trip (which returns ZE_RESULT_ERROR_INVALID_ARGUMENT on current NEO drivers) with a direct file-descriptor approach: - ZeIPCPhysMemHandleData now stores int Fd + size_t Size instead of ze_ipc_mem_handle_t + size_t Size - urIPCGetPhysMemHandleExp: export fd via zePhysicalMemGetProperties and store it directly; the fd lifetime is owned by the handle data - urIPCPutPhysMemHandleExp: close(fd) and delete handle data - urIPCOpenPhysMemHandleExp: use the fd directly in ze_external_memory_import_fd_t without converting via zeMemGetFileDescriptorFromIpcHandleExp - urIPCClosePhysMemHandleExp: zePhysicalMemDestroy + delete handle - Add explicit null-argument checks (INVALID_NULL_HANDLE / INVALID_NULL_POINTER) at the top of all four functions so that the conformance negative tests pass without reaching L0 Signed-off-by: Lukasz Dorau --- .../source/adapters/level_zero/memory.cpp | 70 ++++++----- .../adapters/level_zero/physical_mem.hpp | 6 +- .../source/adapters/level_zero/v2/memory.cpp | 70 ++++++----- .../conformance/virtual_memory/CMakeLists.txt | 4 + .../urIPCClosePhysMemHandleExp.cpp | 25 ++++ .../urIPCGetPhysMemHandleExp.cpp | 53 +++++++++ .../urIPCOpenPhysMemHandleExp.cpp | 57 +++++++++ .../urIPCPhysMemHandleExpFixtures.hpp | 92 +++++++++++++++ .../urIPCPutPhysMemHandleExp.cpp | 25 ++++ .../virtual_memory/urPhysicalMemCreate.cpp | 2 +- .../virtual_memory/ze_ipc_phys_mem_e2e.cpp | 109 ++++++++++++++++++ 11 files changed, 437 insertions(+), 76 deletions(-) create mode 100644 unified-runtime/test/conformance/virtual_memory/urIPCClosePhysMemHandleExp.cpp create mode 100644 unified-runtime/test/conformance/virtual_memory/urIPCGetPhysMemHandleExp.cpp create mode 100644 unified-runtime/test/conformance/virtual_memory/urIPCOpenPhysMemHandleExp.cpp create mode 100644 unified-runtime/test/conformance/virtual_memory/urIPCPhysMemHandleExpFixtures.hpp create mode 100644 unified-runtime/test/conformance/virtual_memory/urIPCPutPhysMemHandleExp.cpp create mode 100644 unified-runtime/test/conformance/virtual_memory/ze_ipc_phys_mem_e2e.cpp diff --git a/unified-runtime/source/adapters/level_zero/memory.cpp b/unified-runtime/source/adapters/level_zero/memory.cpp index 29c54881e58b..ac469988edb6 100644 --- a/unified-runtime/source/adapters/level_zero/memory.cpp +++ b/unified-runtime/source/adapters/level_zero/memory.cpp @@ -2017,12 +2017,14 @@ ur_result_t urIPCGetPhysMemHandleExp(ur_context_handle_t hContext, ur_physical_mem_handle_t hPhysMem, void **ppIPCPhysMemHandleData, size_t *pIPCPhysMemHandleDataSizeRet) { - // Fast path: size-only query. - if (!ppIPCPhysMemHandleData) { - if (pIPCPhysMemHandleDataSizeRet) - *pIPCPhysMemHandleDataSizeRet = sizeof(ZeIPCPhysMemHandleData); - return UR_RESULT_SUCCESS; - } + if (!hContext) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + if (!hPhysMem) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + if (!ppIPCPhysMemHandleData) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + if (!pIPCPhysMemHandleDataSizeRet) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; // Export the physical memory object as an opaque file descriptor. ze_external_memory_export_fd_t ExportFd = {}; @@ -2037,40 +2039,32 @@ ur_result_t urIPCGetPhysMemHandleExp(ur_context_handle_t hContext, ZE2UR_CALL(zePhysicalMemGetProperties, (hContext->getZeHandle(), hPhysMem->ZePhysicalMem, &Props)); - // Convert the file descriptor to an opaque Level Zero IPC handle so that - // the caller can pass it across process boundaries. - ze_ipc_mem_handle_t ZeIpcHandle = {}; - ze_result_t ZeRes = zeMemGetIpcHandleFromFileDescriptorExp( - hContext->getZeHandle(), static_cast(ExportFd.fd), - &ZeIpcHandle); - // The exported fd is no longer needed once the IPC handle is obtained. - close(ExportFd.fd); - if (ZeRes != ZE_RESULT_SUCCESS) - return ze2urResult(ZeRes); - auto *HandleData = new (std::nothrow) ZeIPCPhysMemHandleData; if (!HandleData) { - zeMemPutIpcHandle(hContext->getZeHandle(), ZeIpcHandle); + close(ExportFd.fd); return UR_RESULT_ERROR_OUT_OF_HOST_MEMORY; } - HandleData->ZeHandle = ZeIpcHandle; + HandleData->Fd = ExportFd.fd; HandleData->Size = hPhysMem->Size; *ppIPCPhysMemHandleData = HandleData; - if (pIPCPhysMemHandleDataSizeRet) - *pIPCPhysMemHandleDataSizeRet = sizeof(ZeIPCPhysMemHandleData); + *pIPCPhysMemHandleDataSizeRet = sizeof(ZeIPCPhysMemHandleData); return UR_RESULT_SUCCESS; } ur_result_t urIPCPutPhysMemHandleExp(ur_context_handle_t hContext, void *pIPCPhysMemHandleData) { + if (!hContext) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + if (!pIPCPhysMemHandleData) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + auto *HandleData = static_cast(pIPCPhysMemHandleData); - ur_result_t Res = ze2urResult( - zeMemPutIpcHandle(hContext->getZeHandle(), HandleData->ZeHandle)); + close(HandleData->Fd); delete HandleData; - return Res; + return UR_RESULT_SUCCESS; } ur_result_t urIPCOpenPhysMemHandleExp(ur_context_handle_t hContext, @@ -2078,22 +2072,25 @@ ur_result_t urIPCOpenPhysMemHandleExp(ur_context_handle_t hContext, void *pIPCPhysMemHandleData, size_t ipcPhysMemHandleDataSize, ur_physical_mem_handle_t *phPhysMem) { + if (!hContext) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + if (!hDevice) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + if (!pIPCPhysMemHandleData) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + if (!phPhysMem) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; if (ipcPhysMemHandleDataSize != sizeof(ZeIPCPhysMemHandleData)) return UR_RESULT_ERROR_INVALID_VALUE; auto *HandleData = static_cast(pIPCPhysMemHandleData); - // Convert the opaque IPC handle back to a file descriptor for import. - uint64_t Fd = 0; - ZE2UR_CALL(zeMemGetFileDescriptorFromIpcHandleExp, - (hContext->getZeHandle(), HandleData->ZeHandle, &Fd)); - ze_external_memory_import_fd_t ImportFd = {}; ImportFd.stype = ZE_STRUCTURE_TYPE_EXTERNAL_MEMORY_IMPORT_FD; ImportFd.pNext = nullptr; ImportFd.flags = ZE_EXTERNAL_MEMORY_TYPE_FLAG_OPAQUE_FD; - ImportFd.fd = static_cast(Fd); + ImportFd.fd = HandleData->Fd; ZeStruct PhysMemDesc; PhysMemDesc.pNext = &ImportFd; @@ -2101,12 +2098,8 @@ ur_result_t urIPCOpenPhysMemHandleExp(ur_context_handle_t hContext, PhysMemDesc.size = HandleData->Size; ze_physical_mem_handle_t ZePhysMem; - ze_result_t ZeRes = zePhysicalMemCreate( - hContext->getZeHandle(), hDevice->ZeDevice, &PhysMemDesc, &ZePhysMem); - // The imported fd is no longer needed once zePhysicalMemCreate returns. - close(static_cast(Fd)); - if (ZeRes != ZE_RESULT_SUCCESS) - return ze2urResult(ZeRes); + ZE2UR_CALL(zePhysicalMemCreate, (hContext->getZeHandle(), hDevice->ZeDevice, + &PhysMemDesc, &ZePhysMem)); try { *phPhysMem = new ur_physical_mem_handle_t_(ZePhysMem, hContext, HandleData->Size); @@ -2122,6 +2115,11 @@ ur_result_t urIPCOpenPhysMemHandleExp(ur_context_handle_t hContext, ur_result_t urIPCClosePhysMemHandleExp(ur_context_handle_t hContext, ur_physical_mem_handle_t hPhysMem) { + if (!hContext) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + if (!hPhysMem) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + ur_result_t Res = ze2urResult( zePhysicalMemDestroy(hContext->getZeHandle(), hPhysMem->ZePhysicalMem)); delete hPhysMem; diff --git a/unified-runtime/source/adapters/level_zero/physical_mem.hpp b/unified-runtime/source/adapters/level_zero/physical_mem.hpp index 54e5229d09cc..6bb91f0ce268 100644 --- a/unified-runtime/source/adapters/level_zero/physical_mem.hpp +++ b/unified-runtime/source/adapters/level_zero/physical_mem.hpp @@ -12,10 +12,10 @@ #include "common/ur_ref_count.hpp" // Opaque handle data exchanged between processes for physical memory IPC. -// Contains a Level Zero IPC memory handle and the size of the allocation -// (required by zePhysicalMemCreate on the receiving side). +// Contains the file descriptor exported by zePhysicalMemGetProperties and +// the size of the allocation (required by zePhysicalMemCreate on import). struct ZeIPCPhysMemHandleData { - ze_ipc_mem_handle_t ZeHandle; + int Fd; size_t Size; }; diff --git a/unified-runtime/source/adapters/level_zero/v2/memory.cpp b/unified-runtime/source/adapters/level_zero/v2/memory.cpp index 3541d15a95eb..10c789554b3e 100644 --- a/unified-runtime/source/adapters/level_zero/v2/memory.cpp +++ b/unified-runtime/source/adapters/level_zero/v2/memory.cpp @@ -932,12 +932,14 @@ ur_result_t urIPCGetPhysMemHandleExp(ur_context_handle_t hContext, ur_physical_mem_handle_t hPhysMem, void **ppIPCPhysMemHandleData, size_t *pIPCPhysMemHandleDataSizeRet) { - // Fast path: size-only query. - if (!ppIPCPhysMemHandleData) { - if (pIPCPhysMemHandleDataSizeRet) - *pIPCPhysMemHandleDataSizeRet = sizeof(ZeIPCPhysMemHandleData); - return UR_RESULT_SUCCESS; - } + if (!hContext) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + if (!hPhysMem) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + if (!ppIPCPhysMemHandleData) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + if (!pIPCPhysMemHandleDataSizeRet) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; // Export the physical memory object as an opaque file descriptor. ze_external_memory_export_fd_t ExportFd = {}; @@ -952,40 +954,32 @@ ur_result_t urIPCGetPhysMemHandleExp(ur_context_handle_t hContext, ZE2UR_CALL(zePhysicalMemGetProperties, (hContext->getZeHandle(), hPhysMem->ZePhysicalMem, &Props)); - // Convert the file descriptor to an opaque Level Zero IPC handle so that - // the caller can pass it across process boundaries. - ze_ipc_mem_handle_t ZeIpcHandle = {}; - ze_result_t ZeRes = zeMemGetIpcHandleFromFileDescriptorExp( - hContext->getZeHandle(), static_cast(ExportFd.fd), - &ZeIpcHandle); - // The exported fd is no longer needed once the IPC handle is obtained. - close(ExportFd.fd); - if (ZeRes != ZE_RESULT_SUCCESS) - return ze2urResult(ZeRes); - auto *HandleData = new (std::nothrow) ZeIPCPhysMemHandleData; if (!HandleData) { - zeMemPutIpcHandle(hContext->getZeHandle(), ZeIpcHandle); + close(ExportFd.fd); return UR_RESULT_ERROR_OUT_OF_HOST_MEMORY; } - HandleData->ZeHandle = ZeIpcHandle; + HandleData->Fd = ExportFd.fd; HandleData->Size = hPhysMem->Size; *ppIPCPhysMemHandleData = HandleData; - if (pIPCPhysMemHandleDataSizeRet) - *pIPCPhysMemHandleDataSizeRet = sizeof(ZeIPCPhysMemHandleData); + *pIPCPhysMemHandleDataSizeRet = sizeof(ZeIPCPhysMemHandleData); return UR_RESULT_SUCCESS; } ur_result_t urIPCPutPhysMemHandleExp(ur_context_handle_t hContext, void *pIPCPhysMemHandleData) { + if (!hContext) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + if (!pIPCPhysMemHandleData) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + auto *HandleData = static_cast(pIPCPhysMemHandleData); - ur_result_t Res = ze2urResult( - zeMemPutIpcHandle(hContext->getZeHandle(), HandleData->ZeHandle)); + close(HandleData->Fd); delete HandleData; - return Res; + return UR_RESULT_SUCCESS; } ur_result_t urIPCOpenPhysMemHandleExp(ur_context_handle_t hContext, @@ -993,22 +987,25 @@ ur_result_t urIPCOpenPhysMemHandleExp(ur_context_handle_t hContext, void *pIPCPhysMemHandleData, size_t ipcPhysMemHandleDataSize, ur_physical_mem_handle_t *phPhysMem) { + if (!hContext) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + if (!hDevice) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + if (!pIPCPhysMemHandleData) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + if (!phPhysMem) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; if (ipcPhysMemHandleDataSize != sizeof(ZeIPCPhysMemHandleData)) return UR_RESULT_ERROR_INVALID_VALUE; auto *HandleData = static_cast(pIPCPhysMemHandleData); - // Convert the opaque IPC handle back to a file descriptor for import. - uint64_t Fd = 0; - ZE2UR_CALL(zeMemGetFileDescriptorFromIpcHandleExp, - (hContext->getZeHandle(), HandleData->ZeHandle, &Fd)); - ze_external_memory_import_fd_t ImportFd = {}; ImportFd.stype = ZE_STRUCTURE_TYPE_EXTERNAL_MEMORY_IMPORT_FD; ImportFd.pNext = nullptr; ImportFd.flags = ZE_EXTERNAL_MEMORY_TYPE_FLAG_OPAQUE_FD; - ImportFd.fd = static_cast(Fd); + ImportFd.fd = HandleData->Fd; ZeStruct PhysMemDesc; PhysMemDesc.pNext = &ImportFd; @@ -1016,12 +1013,8 @@ ur_result_t urIPCOpenPhysMemHandleExp(ur_context_handle_t hContext, PhysMemDesc.size = HandleData->Size; ze_physical_mem_handle_t ZePhysMem; - ze_result_t ZeRes = zePhysicalMemCreate( - hContext->getZeHandle(), hDevice->ZeDevice, &PhysMemDesc, &ZePhysMem); - // The imported fd is no longer needed once zePhysicalMemCreate returns. - close(static_cast(Fd)); - if (ZeRes != ZE_RESULT_SUCCESS) - return ze2urResult(ZeRes); + ZE2UR_CALL(zePhysicalMemCreate, (hContext->getZeHandle(), hDevice->ZeDevice, + &PhysMemDesc, &ZePhysMem)); try { *phPhysMem = new ur_physical_mem_handle_t_(ZePhysMem, hContext, HandleData->Size); @@ -1037,6 +1030,11 @@ ur_result_t urIPCOpenPhysMemHandleExp(ur_context_handle_t hContext, ur_result_t urIPCClosePhysMemHandleExp(ur_context_handle_t hContext, ur_physical_mem_handle_t hPhysMem) { + if (!hContext) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + if (!hPhysMem) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + ur_result_t Res = ze2urResult( zePhysicalMemDestroy(hContext->getZeHandle(), hPhysMem->ZePhysicalMem)); delete hPhysMem; diff --git a/unified-runtime/test/conformance/virtual_memory/CMakeLists.txt b/unified-runtime/test/conformance/virtual_memory/CMakeLists.txt index 65c6c03108f1..dc79f146e505 100644 --- a/unified-runtime/test/conformance/virtual_memory/CMakeLists.txt +++ b/unified-runtime/test/conformance/virtual_memory/CMakeLists.txt @@ -14,4 +14,8 @@ add_conformance_devices_test(virtual_memory urVirtualMemReserve.cpp urVirtualMemSetAccess.cpp urVirtualMemUnmap.cpp + urIPCGetPhysMemHandleExp.cpp + urIPCPutPhysMemHandleExp.cpp + urIPCOpenPhysMemHandleExp.cpp + urIPCClosePhysMemHandleExp.cpp ) diff --git a/unified-runtime/test/conformance/virtual_memory/urIPCClosePhysMemHandleExp.cpp b/unified-runtime/test/conformance/virtual_memory/urIPCClosePhysMemHandleExp.cpp new file mode 100644 index 000000000000..4ba5783cdcf0 --- /dev/null +++ b/unified-runtime/test/conformance/virtual_memory/urIPCClosePhysMemHandleExp.cpp @@ -0,0 +1,25 @@ +// 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 "urIPCPhysMemHandleExpFixtures.hpp" + +using urIPCClosePhysMemHandleExpTest = urIPCOpenedPhysMemTest; +UUR_INSTANTIATE_DEVICE_TEST_SUITE(urIPCClosePhysMemHandleExpTest); + +TEST_P(urIPCClosePhysMemHandleExpTest, Success) { + ASSERT_SUCCESS(urIPCClosePhysMemHandleExp(context, opened_physical_mem)); + // Prevent TearDown from attempting a second close. + opened_physical_mem = nullptr; +} + +TEST_P(urIPCClosePhysMemHandleExpTest, InvalidNullHandleContext) { + ASSERT_EQ_RESULT(urIPCClosePhysMemHandleExp(nullptr, opened_physical_mem), + UR_RESULT_ERROR_INVALID_NULL_HANDLE); +} + +TEST_P(urIPCClosePhysMemHandleExpTest, InvalidNullHandlePhysMem) { + ASSERT_EQ_RESULT(urIPCClosePhysMemHandleExp(context, nullptr), + UR_RESULT_ERROR_INVALID_NULL_HANDLE); +} diff --git a/unified-runtime/test/conformance/virtual_memory/urIPCGetPhysMemHandleExp.cpp b/unified-runtime/test/conformance/virtual_memory/urIPCGetPhysMemHandleExp.cpp new file mode 100644 index 000000000000..7d0767972677 --- /dev/null +++ b/unified-runtime/test/conformance/virtual_memory/urIPCGetPhysMemHandleExp.cpp @@ -0,0 +1,53 @@ +// 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 "urIPCPhysMemHandleExpFixtures.hpp" + +using urIPCGetPhysMemHandleExpTest = urIPCPhysMemTest; +UUR_INSTANTIATE_DEVICE_TEST_SUITE(urIPCGetPhysMemHandleExpTest); + +TEST_P(urIPCGetPhysMemHandleExpTest, Success) { + void *ipc_handle_data = nullptr; + size_t ipc_handle_size = 0; + ur_result_t res = urIPCGetPhysMemHandleExp( + context, physical_mem, &ipc_handle_data, &ipc_handle_size); + if (res == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { + GTEST_SKIP() << "IPC physical memory handle export is not supported."; + } + ASSERT_SUCCESS(res); + ASSERT_NE(ipc_handle_data, nullptr); + ASSERT_NE(ipc_handle_size, 0U); + ASSERT_SUCCESS(urIPCPutPhysMemHandleExp(context, ipc_handle_data)); +} + +TEST_P(urIPCGetPhysMemHandleExpTest, InvalidNullHandleContext) { + void *ipc_handle_data = nullptr; + size_t ipc_handle_size = 0; + ASSERT_EQ_RESULT(urIPCGetPhysMemHandleExp(nullptr, physical_mem, + &ipc_handle_data, &ipc_handle_size), + UR_RESULT_ERROR_INVALID_NULL_HANDLE); +} + +TEST_P(urIPCGetPhysMemHandleExpTest, InvalidNullHandlePhysMem) { + void *ipc_handle_data = nullptr; + size_t ipc_handle_size = 0; + ASSERT_EQ_RESULT(urIPCGetPhysMemHandleExp(context, nullptr, &ipc_handle_data, + &ipc_handle_size), + UR_RESULT_ERROR_INVALID_NULL_HANDLE); +} + +TEST_P(urIPCGetPhysMemHandleExpTest, InvalidNullPointerIPCHandleData) { + size_t ipc_handle_size = 0; + ASSERT_EQ_RESULT(urIPCGetPhysMemHandleExp(context, physical_mem, nullptr, + &ipc_handle_size), + UR_RESULT_ERROR_INVALID_NULL_POINTER); +} + +TEST_P(urIPCGetPhysMemHandleExpTest, InvalidNullPointerIPCHandleDataSize) { + void *ipc_handle_data = nullptr; + ASSERT_EQ_RESULT(urIPCGetPhysMemHandleExp(context, physical_mem, + &ipc_handle_data, nullptr), + UR_RESULT_ERROR_INVALID_NULL_POINTER); +} diff --git a/unified-runtime/test/conformance/virtual_memory/urIPCOpenPhysMemHandleExp.cpp b/unified-runtime/test/conformance/virtual_memory/urIPCOpenPhysMemHandleExp.cpp new file mode 100644 index 000000000000..0358f69f1390 --- /dev/null +++ b/unified-runtime/test/conformance/virtual_memory/urIPCOpenPhysMemHandleExp.cpp @@ -0,0 +1,57 @@ +// 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 "urIPCPhysMemHandleExpFixtures.hpp" + +using urIPCOpenPhysMemHandleExpTest = urIPCPhysMemHandleTest; +UUR_INSTANTIATE_DEVICE_TEST_SUITE(urIPCOpenPhysMemHandleExpTest); + +TEST_P(urIPCOpenPhysMemHandleExpTest, Success) { + ur_physical_mem_handle_t opened_physical_mem = nullptr; + ASSERT_SUCCESS(urIPCOpenPhysMemHandleExp( + context, device, ipc_handle_data, ipc_handle_size, &opened_physical_mem)); + ASSERT_NE(opened_physical_mem, nullptr); + ASSERT_SUCCESS(urIPCClosePhysMemHandleExp(context, opened_physical_mem)); +} + +TEST_P(urIPCOpenPhysMemHandleExpTest, InvalidNullHandleContext) { + ur_physical_mem_handle_t opened_physical_mem = nullptr; + ASSERT_EQ_RESULT(urIPCOpenPhysMemHandleExp(nullptr, device, ipc_handle_data, + ipc_handle_size, + &opened_physical_mem), + UR_RESULT_ERROR_INVALID_NULL_HANDLE); +} + +TEST_P(urIPCOpenPhysMemHandleExpTest, InvalidNullHandleDevice) { + ur_physical_mem_handle_t opened_physical_mem = nullptr; + ASSERT_EQ_RESULT(urIPCOpenPhysMemHandleExp(context, nullptr, ipc_handle_data, + ipc_handle_size, + &opened_physical_mem), + UR_RESULT_ERROR_INVALID_NULL_HANDLE); +} + +TEST_P(urIPCOpenPhysMemHandleExpTest, InvalidNullPointerIPCHandleData) { + ur_physical_mem_handle_t opened_physical_mem = nullptr; + ASSERT_EQ_RESULT(urIPCOpenPhysMemHandleExp(context, device, nullptr, + ipc_handle_size, + &opened_physical_mem), + UR_RESULT_ERROR_INVALID_NULL_POINTER); +} + +TEST_P(urIPCOpenPhysMemHandleExpTest, InvalidNullPointerPhysMem) { + ASSERT_EQ_RESULT(urIPCOpenPhysMemHandleExp(context, device, ipc_handle_data, + ipc_handle_size, nullptr), + UR_RESULT_ERROR_INVALID_NULL_POINTER); +} + +TEST_P(urIPCOpenPhysMemHandleExpTest, InvalidValue) { + ur_physical_mem_handle_t opened_physical_mem = nullptr; + // Pass a size that differs from the real handle data size to trigger + // UR_RESULT_ERROR_INVALID_VALUE. + ASSERT_EQ_RESULT(urIPCOpenPhysMemHandleExp(context, device, ipc_handle_data, + ipc_handle_size + 1, + &opened_physical_mem), + UR_RESULT_ERROR_INVALID_VALUE); +} diff --git a/unified-runtime/test/conformance/virtual_memory/urIPCPhysMemHandleExpFixtures.hpp b/unified-runtime/test/conformance/virtual_memory/urIPCPhysMemHandleExpFixtures.hpp new file mode 100644 index 000000000000..4134ec097d49 --- /dev/null +++ b/unified-runtime/test/conformance/virtual_memory/urIPCPhysMemHandleExpFixtures.hpp @@ -0,0 +1,92 @@ +// 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 + +// Fixture that creates a physical memory object with the IPC-export flag. +struct urIPCPhysMemTest : uur::urVirtualMemGranularityTest { + void SetUp() override { + UUR_RETURN_ON_FATAL_FAILURE(urVirtualMemGranularityTest::SetUp()); + size = granularity * 256; + + ur_bool_t ipc_support = false; + ASSERT_SUCCESS(urDeviceGetInfo(device, + UR_DEVICE_INFO_IPC_MEMORY_SUPPORT_EXP, + sizeof(ur_bool_t), &ipc_support, nullptr)); + if (!ipc_support) { + GTEST_SKIP() << "IPC memory is not supported."; + } + + ur_physical_mem_properties_t properties{ + UR_STRUCTURE_TYPE_PHYSICAL_MEM_PROPERTIES, nullptr, + UR_PHYSICAL_MEM_FLAG_ENABLE_IPC}; + ur_result_t res = + urPhysicalMemCreate(context, device, size, &properties, &physical_mem); + if (res == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { + GTEST_SKIP() << "IPC physical memory is not supported."; + } + ASSERT_SUCCESS(res); + ASSERT_NE(physical_mem, nullptr); + } + + void TearDown() override { + if (physical_mem) { + EXPECT_SUCCESS(urPhysicalMemRelease(physical_mem)); + } + uur::urVirtualMemGranularityTest::TearDown(); + } + + size_t size = 0; + ur_physical_mem_handle_t physical_mem = nullptr; +}; + +// Fixture that also acquires an IPC handle for the physical memory object. +struct urIPCPhysMemHandleTest : urIPCPhysMemTest { + void SetUp() override { + UUR_RETURN_ON_FATAL_FAILURE(urIPCPhysMemTest::SetUp()); + ur_result_t res = urIPCGetPhysMemHandleExp( + context, physical_mem, &ipc_handle_data, &ipc_handle_size); + if (res == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { + GTEST_SKIP() << "IPC physical memory handle export is not supported."; + } + ASSERT_SUCCESS(res); + ASSERT_NE(ipc_handle_data, nullptr); + ASSERT_NE(ipc_handle_size, 0U); + } + + void TearDown() override { + if (ipc_handle_data) { + EXPECT_SUCCESS(urIPCPutPhysMemHandleExp(context, ipc_handle_data)); + ipc_handle_data = nullptr; + } + urIPCPhysMemTest::TearDown(); + } + + void *ipc_handle_data = nullptr; + size_t ipc_handle_size = 0; +}; + +// Fixture that also opens the IPC handle to produce a second physical_mem. +struct urIPCOpenedPhysMemTest : urIPCPhysMemHandleTest { + void SetUp() override { + UUR_RETURN_ON_FATAL_FAILURE(urIPCPhysMemHandleTest::SetUp()); + ASSERT_SUCCESS(urIPCOpenPhysMemHandleExp(context, device, ipc_handle_data, + ipc_handle_size, + &opened_physical_mem)); + ASSERT_NE(opened_physical_mem, nullptr); + } + + void TearDown() override { + if (opened_physical_mem) { + EXPECT_SUCCESS(urIPCClosePhysMemHandleExp(context, opened_physical_mem)); + opened_physical_mem = nullptr; + } + urIPCPhysMemHandleTest::TearDown(); + } + + ur_physical_mem_handle_t opened_physical_mem = nullptr; +}; diff --git a/unified-runtime/test/conformance/virtual_memory/urIPCPutPhysMemHandleExp.cpp b/unified-runtime/test/conformance/virtual_memory/urIPCPutPhysMemHandleExp.cpp new file mode 100644 index 000000000000..7a824b7ab6ef --- /dev/null +++ b/unified-runtime/test/conformance/virtual_memory/urIPCPutPhysMemHandleExp.cpp @@ -0,0 +1,25 @@ +// 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 "urIPCPhysMemHandleExpFixtures.hpp" + +using urIPCPutPhysMemHandleExpTest = urIPCPhysMemHandleTest; +UUR_INSTANTIATE_DEVICE_TEST_SUITE(urIPCPutPhysMemHandleExpTest); + +TEST_P(urIPCPutPhysMemHandleExpTest, Success) { + ASSERT_SUCCESS(urIPCPutPhysMemHandleExp(context, ipc_handle_data)); + // Prevent TearDown from attempting a second put. + ipc_handle_data = nullptr; +} + +TEST_P(urIPCPutPhysMemHandleExpTest, InvalidNullHandleContext) { + ASSERT_EQ_RESULT(urIPCPutPhysMemHandleExp(nullptr, ipc_handle_data), + UR_RESULT_ERROR_INVALID_NULL_HANDLE); +} + +TEST_P(urIPCPutPhysMemHandleExpTest, InvalidNullPointerIPCHandleData) { + ASSERT_EQ_RESULT(urIPCPutPhysMemHandleExp(context, nullptr), + UR_RESULT_ERROR_INVALID_NULL_POINTER); +} diff --git a/unified-runtime/test/conformance/virtual_memory/urPhysicalMemCreate.cpp b/unified-runtime/test/conformance/virtual_memory/urPhysicalMemCreate.cpp index f7c9538d7929..77ef56462f14 100644 --- a/unified-runtime/test/conformance/virtual_memory/urPhysicalMemCreate.cpp +++ b/unified-runtime/test/conformance/virtual_memory/urPhysicalMemCreate.cpp @@ -55,7 +55,7 @@ using urPhysicalMemCreateWithFlagsParamTest = uur::urPhysicalMemTestWithParam; UUR_DEVICE_TEST_SUITE_WITH_PARAM( urPhysicalMemCreateWithFlagsParamTest, - ::testing::Values(UR_PHYSICAL_MEM_FLAG_TBD), + ::testing::Values(UR_PHYSICAL_MEM_FLAG_ENABLE_IPC), uur::deviceTestWithParamPrinter); TEST_P(urPhysicalMemCreateWithFlagsParamTest, Success) { diff --git a/unified-runtime/test/conformance/virtual_memory/ze_ipc_phys_mem_e2e.cpp b/unified-runtime/test/conformance/virtual_memory/ze_ipc_phys_mem_e2e.cpp new file mode 100644 index 000000000000..c9b167f53e48 --- /dev/null +++ b/unified-runtime/test/conformance/virtual_memory/ze_ipc_phys_mem_e2e.cpp @@ -0,0 +1,109 @@ +// 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 +// +// End-to-end diagnostic test for the Level Zero IPC physical memory APIs. +// Exercises the full export → put → open → close flow using L0 directly. +// +// Build (requires NEO ze_loader): +// g++ -std=c++17 \ +// -I/home/ldorau/work/neo_workspace/level_zero/include \ +// -L/home/ldorau/work/neo_workspace/level_zero/lib \ +// -Wl,-rpath,/home/ldorau/work/neo_workspace/level_zero/lib \ +// -l:libze_loader.so.1 \ +// ze_ipc_phys_mem_e2e.cpp -o ze_ipc_phys_mem_e2e +// +// Run (with NEO driver): +// L0 vNEO Release ./ze_ipc_phys_mem_e2e + +#include +#include +#include + +#define CHECK(call) \ + do { \ + ze_result_t r = (call); \ + if (r != ZE_RESULT_SUCCESS) { \ + printf("FAILED %s = 0x%x\n", #call, r); \ + return 1; \ + } \ + printf("OK %s\n", #call); \ + } while (0) + +int main() { + CHECK(zeInit(0)); + + ze_driver_handle_t driver; + uint32_t n = 1; + CHECK(zeDriverGet(&n, &driver)); + + ze_device_handle_t device; + n = 1; + CHECK(zeDeviceGet(driver, &n, &device)); + + ze_context_desc_t ctxDesc{ZE_STRUCTURE_TYPE_CONTEXT_DESC, nullptr, 0}; + ze_context_handle_t ctx; + CHECK(zeContextCreate(driver, &ctxDesc, &ctx)); + + // Create physical memory with export descriptor so it can be shared via IPC. + ze_external_memory_export_desc_t exportDesc{}; + exportDesc.stype = ZE_STRUCTURE_TYPE_EXTERNAL_MEMORY_EXPORT_DESC; + exportDesc.flags = ZE_EXTERNAL_MEMORY_TYPE_FLAG_OPAQUE_FD; + + ze_physical_mem_desc_t desc{}; + desc.stype = ZE_STRUCTURE_TYPE_PHYSICAL_MEM_DESC; + desc.pNext = &exportDesc; + desc.size = 2 * 1024 * 1024; + + ze_physical_mem_handle_t physMem; + CHECK(zePhysicalMemCreate(ctx, device, &desc, &physMem)); + + // Export the physical memory object as an opaque file descriptor. + ze_external_memory_export_fd_t exportFd{}; + exportFd.stype = ZE_STRUCTURE_TYPE_EXTERNAL_MEMORY_EXPORT_FD; + exportFd.flags = ZE_EXTERNAL_MEMORY_TYPE_FLAG_OPAQUE_FD; + + ze_physical_mem_properties_t props{}; + props.stype = ZE_STRUCTURE_TYPE_PHYSICAL_MEM_PROPERTIES; + props.pNext = &exportFd; + CHECK(zePhysicalMemGetProperties(ctx, physMem, &props)); + printf(" Exported fd = %d\n", exportFd.fd); + + // Convert the file descriptor to an opaque IPC handle. + ze_ipc_mem_handle_t ipcHandle{}; + CHECK(zeMemGetIpcHandleFromFileDescriptorExp(ctx, (uint64_t)exportFd.fd, + &ipcHandle)); + close(exportFd.fd); + printf(" Got IPC handle\n"); + + // Simulate the receiving side: convert IPC handle back to a file descriptor. + uint64_t fd2 = 0; + CHECK(zeMemGetFileDescriptorFromIpcHandleExp(ctx, ipcHandle, &fd2)); + printf(" Received fd2 = %lu\n", fd2); + + // Open the imported physical memory object from the received fd. + ze_external_memory_import_fd_t importFd{}; + importFd.stype = ZE_STRUCTURE_TYPE_EXTERNAL_MEMORY_IMPORT_FD; + importFd.flags = ZE_EXTERNAL_MEMORY_TYPE_FLAG_OPAQUE_FD; + importFd.fd = (int)fd2; + + ze_physical_mem_desc_t importDesc{}; + importDesc.stype = ZE_STRUCTURE_TYPE_PHYSICAL_MEM_DESC; + importDesc.pNext = &importFd; + importDesc.size = 2 * 1024 * 1024; + + ze_physical_mem_handle_t importedMem; + CHECK(zePhysicalMemCreate(ctx, device, &importDesc, &importedMem)); + close((int)fd2); + printf(" Opened imported physical mem\n"); + + // Clean up in reverse order. + CHECK(zePhysicalMemDestroy(ctx, importedMem)); + CHECK(zeMemPutIpcHandle(ctx, ipcHandle)); + CHECK(zePhysicalMemDestroy(ctx, physMem)); + CHECK(zeContextDestroy(ctx)); + + printf("\nAll OK!\n"); + return 0; +} From 847ea708fb776d622fe58a783a41f08279ab4850 Mon Sep 17 00:00:00 2001 From: Lukasz Dorau Date: Fri, 22 May 2026 09:08:13 +0000 Subject: [PATCH 09/17] [UR][L0] Add cross-process IPC support for physical_mem handles MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Use pidfd_getfd(2) to transfer DMA-BUF file descriptors between processes, enabling true cross-process sharing of physical_mem objects via the IPC API. Changes: - ZeIPCPhysMemHandleData: add Pid field (exporting process PID) - urIPCGetPhysMemHandleExp: store getpid() in handle data; keep fd open until urIPCPutPhysMemHandleExp closes it - urIPCOpenPhysMemHandleExp: same-process → dup(), cross-process → pidfd_open(Pid) + pidfd_getfd() to obtain a local copy of the DMA-BUF fd; close the local copy after zePhysicalMemCreate (fixes the fd leak bug identified in code review) - Add UR_DEVICE_INFO_IPC_PHYSICAL_MEMORY_SUPPORT_EXP = 0x2024 to virtual_memory.yml spec and regenerate ur_api.h / ur_print.hpp - Implement new device info in L0 adapter (returns true on Linux) - Add false stub in CUDA, HIP, OpenCL, native_cpu adapters Print step-by-step status for both spawner and consumer so users can follow what is happening and clearly see PASSED/FAILED outcome. Signed-off-by: Lukasz Dorau --- sycl/source/detail/device_impl.hpp | 5 + .../detail/ur_device_info_ret_types.inc | 1 + .../Experimental/ipc_physical_memory.cpp | 236 ++++++++++++++++++ .../include/unified-runtime/ur_api.h | 3 + .../include/unified-runtime/ur_print.hpp | 16 ++ .../scripts/core/virtual_memory.yml | 11 + .../source/adapters/cuda/device.cpp | 2 + .../source/adapters/hip/device.cpp | 2 + .../source/adapters/level_zero/device.cpp | 6 + .../source/adapters/level_zero/memory.cpp | 37 ++- .../adapters/level_zero/physical_mem.hpp | 9 +- .../source/adapters/level_zero/v2/memory.cpp | 37 ++- .../source/adapters/native_cpu/device.cpp | 2 + .../source/adapters/opencl/device.cpp | 2 + 14 files changed, 360 insertions(+), 9 deletions(-) create mode 100644 sycl/test-e2e/Experimental/ipc_physical_memory.cpp diff --git a/sycl/source/detail/device_impl.hpp b/sycl/source/detail/device_impl.hpp index f2d461c007ea..09e37bca1b7a 100644 --- a/sycl/source/detail/device_impl.hpp +++ b/sycl/source/detail/device_impl.hpp @@ -1608,6 +1608,11 @@ class device_impl { return get_info_impl_nocheck() .value_or(0); } + CASE(ext_oneapi_ipc_physical_memory) { + return get_info_impl_nocheck< + UR_DEVICE_INFO_IPC_PHYSICAL_MEMORY_SUPPORT_EXP>() + .value_or(0); + } CASE(ext_oneapi_device_wait) { return get_info_impl_nocheck() .value_or(0); diff --git a/sycl/source/detail/ur_device_info_ret_types.inc b/sycl/source/detail/ur_device_info_ret_types.inc index 9a51a0ed99ee..e7e4ae7862c8 100644 --- a/sycl/source/detail/ur_device_info_ret_types.inc +++ b/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_PHYSICAL_MEMORY_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) diff --git a/sycl/test-e2e/Experimental/ipc_physical_memory.cpp b/sycl/test-e2e/Experimental/ipc_physical_memory.cpp new file mode 100644 index 000000000000..5ba8132a5042 --- /dev/null +++ b/sycl/test-e2e/Experimental/ipc_physical_memory.cpp @@ -0,0 +1,236 @@ +// REQUIRES: aspect-ext_oneapi_virtual_mem && aspect-ext_oneapi_ipc_physical_memory + +// UNSUPPORTED: level_zero && windows +// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/0000 + +// DEFINE: %{cpp20} = %if cl_options %{/clang:-std=c++20%} %else %{-std=c++20%} + +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out +// RUN: %{build} -DUSE_VIEW %{cpp20} -o %t.view.out +// RUN: %{run} %t.view.out + +// Tests cross-process IPC for physical_mem objects: the spawner creates a +// physical_mem with enable_ipc, maps it, writes data, serializes the IPC +// handle to a file, and spawns a consumer. The consumer opens the physical +// mem handle, verifies the data, writes new data, then exits. The spawner +// verifies the consumer's writes. +// +// Cross-process fd sharing uses pidfd_getfd(2) (Linux 5.6+) internally. +// The spawner must make itself ptrace-accessible so the consumer can copy +// the DMA-BUF fd via pidfd_getfd. + +#include +#include +#include +#include + +#include +#include +#include +#include + +#if defined(__linux__) +#include +#include +#include +#endif // defined(__linux__) + +namespace syclexp = sycl::ext::oneapi::experimental; +namespace syclipc = sycl::ext::oneapi::experimental::ipc; + +constexpr size_t N = 32; +constexpr const char *CommsFile = "ipc_phys_mem_comms.txt"; + +// Return the smallest multiple of Granularity that is >= Bytes. +static size_t alignUp(size_t Bytes, size_t Granularity) { + return ((Bytes + Granularity - 1) / Granularity) * Granularity; +} + +// Return the granularity to use for both physical and virtual allocations. +static size_t getGranularity(const sycl::device &Dev, + const sycl::context &Ctx) { + size_t CtxGran = syclexp::get_mem_granularity(Ctx); + size_t DevGran = syclexp::get_mem_granularity(Dev, Ctx); + // Use the LCM so the size is aligned to both constraints. + size_t GCD = CtxGran; + size_t Rem = DevGran % GCD; + while (Rem != 0) { + std::swap(GCD, Rem); + Rem %= GCD; + } + return (DevGran / GCD) * CtxGran; +} + +int spawner(int argc, char *argv[]) { + assert(argc == 1); + sycl::queue Q; + sycl::context Ctx = Q.get_context(); + sycl::device Dev = Q.get_device(); + + std::cout << "[Spawner] Device: " << Dev.get_info() + << "\n"; + std::cout << "[Spawner] Creating physical_mem (" << N + << " ints) with IPC support...\n"; + +#if defined(__linux__) + // Allow any process (the spawned consumer) to copy our DMA-BUF fd via + // pidfd_getfd(2), which internally uses PTRACE_MODE_ATTACH permissions. + // Without this, pidfd_getfd fails with EPERM under ptrace_scope=1. + prctl(PR_SET_PTRACER, PR_SET_PTRACER_ANY); +#endif // defined(__linux__) + + const size_t AlignedByteSize = + alignUp(N * sizeof(int), getGranularity(Dev, Ctx)); + + // Create a physical_mem that can be shared via IPC. + syclexp::physical_mem PhysMem{Dev, Ctx, AlignedByteSize, + syclexp::properties{syclexp::enable_ipc{}}}; + + // Reserve virtual address space and map the physical memory into it. + std::cout << "[Spawner] Mapping " << AlignedByteSize + << " bytes of virtual address space...\n"; + uintptr_t VAddr = syclexp::reserve_virtual_mem(AlignedByteSize, Ctx); + int *DataPtr = reinterpret_cast(PhysMem.map( + VAddr, AlignedByteSize, syclexp::address_access_mode::read_write)); + + // Initialize: write [0, 1, ..., N-1] into the mapped memory. + std::cout << "[Spawner] Writing initial data [0.." << N - 1 << "]...\n"; + Q.parallel_for(N, [=](sycl::item<1> I) { + DataPtr[I] = static_cast(I.get_linear_id()); + }).wait(); + + { + // Obtain the IPC handle and serialize it together with AlignedByteSize. + std::cout << "[Spawner] Exporting IPC handle to '" << CommsFile << "'...\n"; + syclipc::handle Handle = syclipc::physical_memory::get(PhysMem); + +#ifdef USE_VIEW + syclipc::handle_data_view_t HandleData = Handle.data_view(); +#else + syclipc::handle_data_t HandleData = Handle.data(); +#endif + + size_t HandleDataSize = HandleData.size(); + std::fstream FS(CommsFile, std::ios_base::out | std::ios_base::binary); + FS.write(reinterpret_cast(&AlignedByteSize), sizeof(size_t)); + FS.write(reinterpret_cast(&HandleDataSize), sizeof(size_t)); + FS.write(reinterpret_cast(HandleData.data()), HandleDataSize); + FS.close(); + + // Spawn the consumer process; it reads the comms file, opens the + // physical_mem, verifies the spawner's data, and writes new values. + std::string Cmd = std::string{argv[0]} + " 1"; + std::cout << "[Spawner] Spawning consumer: " << Cmd << "\n"; + std::system(Cmd.c_str()); + + // Release the IPC handle (closes the exported fd). + syclipc::physical_memory::put(Handle, Ctx); + } + + // After consumer exits, verify it wrote [N, N-1, ..., 1]. + std::cout << "[Spawner] Verifying consumer wrote [" << N << ".." << 1 + << "]...\n"; + int Failures = 0; + int Read[N] = {0}; + Q.copy(DataPtr, Read, N).wait(); + for (size_t I = 0; I < N; ++I) { + const int Expected = static_cast(N - I); + if (Read[I] != Expected) { + ++Failures; + std::cout << "[Spawner] MISMATCH at [" << I << "]: got " << Read[I] + << ", expected " << Expected << "\n"; + } + } + + // Cleanup. + syclexp::unmap(DataPtr, AlignedByteSize, Ctx); + syclexp::free_virtual_mem(VAddr, AlignedByteSize, Ctx); + + if (Failures == 0) + std::cout << "[Spawner] PASSED\n"; + else + std::cout << "[Spawner] FAILED (" << Failures << " mismatches)\n"; + return Failures; +} + +int consumer() { + sycl::queue Q; + sycl::context Ctx = Q.get_context(); + sycl::device Dev = Q.get_device(); + + std::cout << "[Consumer] Device: " << Dev.get_info() + << "\n"; + std::cout << "[Consumer] Reading IPC handle from '" << CommsFile << "'...\n"; + + // Read the serialized handle from the comms file. + std::fstream FS(CommsFile, std::ios_base::in | std::ios_base::binary); + size_t AlignedByteSize = 0; + FS.read(reinterpret_cast(&AlignedByteSize), sizeof(size_t)); + size_t HandleSize = 0; + FS.read(reinterpret_cast(&HandleSize), sizeof(size_t)); + std::unique_ptr HandleBytes{new std::byte[HandleSize]}; + FS.read(reinterpret_cast(HandleBytes.get()), HandleSize); + FS.close(); + + // Open the physical_mem from the IPC handle. The resulting object will + // call urIPCClosePhysMemHandleExp when destroyed. + std::cout << "[Consumer] Opening physical_mem from IPC handle (" << HandleSize + << " bytes)...\n"; +#ifdef USE_VIEW + syclipc::handle_data_view_t HandleData{HandleBytes.get(), HandleSize}; +#else + syclipc::handle_data_t HandleData{HandleBytes.get(), + HandleBytes.get() + HandleSize}; +#endif + + syclexp::physical_mem PhysMem = + syclipc::physical_memory::open(HandleData, Ctx, Dev); + + // Map the opened physical_mem at a fresh virtual address. + std::cout << "[Consumer] Mapping " << AlignedByteSize + << " bytes of virtual address space...\n"; + uintptr_t VAddr = syclexp::reserve_virtual_mem(AlignedByteSize, Ctx); + int *DataPtr = reinterpret_cast(PhysMem.map( + VAddr, AlignedByteSize, syclexp::address_access_mode::read_write)); + + // Verify the spawner wrote [0, 1, ..., N-1]. + std::cout << "[Consumer] Verifying spawner wrote [0.." << N - 1 << "]...\n"; + int Failures = 0; + int Read[N] = {0}; + Q.copy(DataPtr, Read, N).wait(); + for (size_t I = 0; I < N; ++I) { + const int Expected = static_cast(I); + if (Read[I] != Expected) { + ++Failures; + std::cout << "[Consumer] MISMATCH at [" << I << "]: got " << Read[I] + << ", expected " << Expected << "\n"; + } + } + + // Write [N, N-1, ..., 1] so the spawner can verify cross-process writes. + std::cout << "[Consumer] Writing new data [" << N << ".." << 1 << "]...\n"; + Q.parallel_for(N, [=](sycl::item<1> I) { + DataPtr[I] = static_cast(N - I.get_linear_id()); + }).wait(); + + // Cleanup virtual-address resources; PhysMem destructor closes the handle. + syclexp::unmap(DataPtr, AlignedByteSize, Ctx); + syclexp::free_virtual_mem(VAddr, AlignedByteSize, Ctx); + + if (Failures == 0) + std::cout << "[Consumer] PASSED\n"; + else + std::cout << "[Consumer] FAILED (" << Failures << " mismatches)\n"; + return Failures; +} + +int main(int argc, char *argv[]) { + if (argc == 1) + std::cout << "=== ipc_physical_memory test" +#ifdef USE_VIEW + << " (USE_VIEW)" +#endif + << " ===\n"; + return argc == 1 ? spawner(argc, argv) : consumer(); +} diff --git a/unified-runtime/include/unified-runtime/ur_api.h b/unified-runtime/include/unified-runtime/ur_api.h index 6c252e35a4c9..e18d7000cf2e 100644 --- a/unified-runtime/include/unified-runtime/ur_api.h +++ b/unified-runtime/include/unified-runtime/ur_api.h @@ -2511,6 +2511,9 @@ typedef enum ur_device_info_t { /// [::ur_bool_t] returns true if the device supports inter-process /// communicable memory handles UR_DEVICE_INFO_IPC_MEMORY_SUPPORT_EXP = 0x2023, + /// [::ur_bool_t] returns true if the device supports inter-process + /// communicable physical memory handles + UR_DEVICE_INFO_IPC_PHYSICAL_MEMORY_SUPPORT_EXP = 0x2024, /// [::ur_bool_t] returns true if the device supports enqueueing of /// allocations and frees. UR_DEVICE_INFO_ASYNC_USM_ALLOCATIONS_SUPPORT_EXP = 0x2050, diff --git a/unified-runtime/include/unified-runtime/ur_print.hpp b/unified-runtime/include/unified-runtime/ur_print.hpp index 50a876b67ca6..16ab1442abb6 100644 --- a/unified-runtime/include/unified-runtime/ur_print.hpp +++ b/unified-runtime/include/unified-runtime/ur_print.hpp @@ -3318,6 +3318,9 @@ inline std::ostream &operator<<(std::ostream &os, enum ur_device_info_t value) { case UR_DEVICE_INFO_IPC_MEMORY_SUPPORT_EXP: os << "UR_DEVICE_INFO_IPC_MEMORY_SUPPORT_EXP"; break; + case UR_DEVICE_INFO_IPC_PHYSICAL_MEMORY_SUPPORT_EXP: + os << "UR_DEVICE_INFO_IPC_PHYSICAL_MEMORY_SUPPORT_EXP"; + break; case UR_DEVICE_INFO_ASYNC_USM_ALLOCATIONS_SUPPORT_EXP: os << "UR_DEVICE_INFO_ASYNC_USM_ALLOCATIONS_SUPPORT_EXP"; break; @@ -5603,6 +5606,19 @@ inline ur_result_t printTagged(std::ostream &os, const void *ptr, os << ")"; } break; + case UR_DEVICE_INFO_IPC_PHYSICAL_MEMORY_SUPPORT_EXP: { + const ur_bool_t *tptr = (const ur_bool_t *)ptr; + if (sizeof(ur_bool_t) > size) { + os << "invalid size (is: " << size + << ", expected: >=" << sizeof(ur_bool_t) << ")"; + return UR_RESULT_ERROR_INVALID_SIZE; + } + os << (const void *)(tptr) << " ("; + + os << *tptr; + + os << ")"; + } break; case UR_DEVICE_INFO_ASYNC_USM_ALLOCATIONS_SUPPORT_EXP: { const ur_bool_t *tptr = (const ur_bool_t *)ptr; if (sizeof(ur_bool_t) > size) { diff --git a/unified-runtime/scripts/core/virtual_memory.yml b/unified-runtime/scripts/core/virtual_memory.yml index 4fe51b74dd55..d1b841099486 100644 --- a/unified-runtime/scripts/core/virtual_memory.yml +++ b/unified-runtime/scripts/core/virtual_memory.yml @@ -11,6 +11,17 @@ type: header desc: "Intel $OneApi Unified Runtime APIs" ordinal: "4" +--- #-------------------------------------------------------------------------- +type: enum +extend: true +typed_etors: true +desc: "Extension enums to $x_device_info_t to support inter-process communicable physical memory handles." +name: $x_device_info_t +etors: + - name: IPC_PHYSICAL_MEMORY_SUPPORT_EXP + value: "0x2024" + desc: "[$x_bool_t] returns true if the device supports inter-process communicable physical memory handles" + --- #-------------------------------------------------------------------------- type: enum desc: "Virtual memory granularity info" diff --git a/unified-runtime/source/adapters/cuda/device.cpp b/unified-runtime/source/adapters/cuda/device.cpp index d213d51eaeb9..9ef56d297e33 100644 --- a/unified-runtime/source/adapters/cuda/device.cpp +++ b/unified-runtime/source/adapters/cuda/device.cpp @@ -1173,6 +1173,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, hDevice->get())); return ReturnValue(static_cast(IPCSupported)); } + case UR_DEVICE_INFO_IPC_PHYSICAL_MEMORY_SUPPORT_EXP: + return ReturnValue(false); case UR_DEVICE_INFO_COMMAND_BUFFER_SUPPORT_EXP: case UR_DEVICE_INFO_COMMAND_BUFFER_EVENT_SUPPORT_EXP: return ReturnValue(true); diff --git a/unified-runtime/source/adapters/hip/device.cpp b/unified-runtime/source/adapters/hip/device.cpp index f07a31c03ce4..e43f7e58164f 100644 --- a/unified-runtime/source/adapters/hip/device.cpp +++ b/unified-runtime/source/adapters/hip/device.cpp @@ -1021,6 +1021,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, static_cast(0)); case UR_DEVICE_INFO_IPC_MEMORY_SUPPORT_EXP: return ReturnValue(false); + case UR_DEVICE_INFO_IPC_PHYSICAL_MEMORY_SUPPORT_EXP: + return ReturnValue(false); case UR_DEVICE_INFO_COMMAND_BUFFER_SUPPORT_EXP: { int RuntimeVersion = 0; UR_CHECK_ERROR(hipRuntimeGetVersion(&RuntimeVersion)); diff --git a/unified-runtime/source/adapters/level_zero/device.cpp b/unified-runtime/source/adapters/level_zero/device.cpp index 03b570bda31d..9af27c402efa 100644 --- a/unified-runtime/source/adapters/level_zero/device.cpp +++ b/unified-runtime/source/adapters/level_zero/device.cpp @@ -1379,6 +1379,12 @@ ur_result_t urDeviceGetInfo( return ReturnValue(false); #else return ReturnValue(true); +#endif + case UR_DEVICE_INFO_IPC_PHYSICAL_MEMORY_SUPPORT_EXP: +#ifdef _WIN32 + return ReturnValue(false); +#else + return ReturnValue(true); #endif case UR_DEVICE_INFO_ASYNC_BARRIER: return ReturnValue(false); diff --git a/unified-runtime/source/adapters/level_zero/memory.cpp b/unified-runtime/source/adapters/level_zero/memory.cpp index ac469988edb6..473455f08b02 100644 --- a/unified-runtime/source/adapters/level_zero/memory.cpp +++ b/unified-runtime/source/adapters/level_zero/memory.cpp @@ -10,6 +10,7 @@ #include #include #include +#include #include #include @@ -2045,6 +2046,10 @@ ur_result_t urIPCGetPhysMemHandleExp(ur_context_handle_t hContext, return UR_RESULT_ERROR_OUT_OF_HOST_MEMORY; } + // Store the exporting process's PID and fd. The fd stays open until + // urIPCPutPhysMemHandleExp is called. Cross-process consumers use + // pidfd_getfd(2) to obtain their own duplicate of this fd. + HandleData->Pid = getpid(); HandleData->Fd = ExportFd.fd; HandleData->Size = hPhysMem->Size; @@ -2086,11 +2091,31 @@ ur_result_t urIPCOpenPhysMemHandleExp(ur_context_handle_t hContext, auto *HandleData = static_cast(pIPCPhysMemHandleData); + // Obtain a usable fd in the current process. For same-process opens + // (e.g. conformance tests) dup() suffices. For cross-process opens + // use pidfd_getfd(2) which requires the exporting process to be + // ptrace-accessible (e.g. via prctl(PR_SET_PTRACER, PR_SET_PTRACER_ANY)). + int ImportFdNum = -1; + if (HandleData->Pid == getpid()) { + ImportFdNum = dup(HandleData->Fd); + if (ImportFdNum < 0) + return UR_RESULT_ERROR_INVALID_VALUE; + } else { + int PidFd = static_cast(syscall(SYS_pidfd_open, HandleData->Pid, 0)); + if (PidFd < 0) + return UR_RESULT_ERROR_INVALID_VALUE; + ImportFdNum = + static_cast(syscall(SYS_pidfd_getfd, PidFd, HandleData->Fd, 0)); + close(PidFd); + if (ImportFdNum < 0) + return UR_RESULT_ERROR_INVALID_VALUE; + } + ze_external_memory_import_fd_t ImportFd = {}; ImportFd.stype = ZE_STRUCTURE_TYPE_EXTERNAL_MEMORY_IMPORT_FD; ImportFd.pNext = nullptr; ImportFd.flags = ZE_EXTERNAL_MEMORY_TYPE_FLAG_OPAQUE_FD; - ImportFd.fd = HandleData->Fd; + ImportFd.fd = ImportFdNum; ZeStruct PhysMemDesc; PhysMemDesc.pNext = &ImportFd; @@ -2098,8 +2123,14 @@ ur_result_t urIPCOpenPhysMemHandleExp(ur_context_handle_t hContext, PhysMemDesc.size = HandleData->Size; ze_physical_mem_handle_t ZePhysMem; - ZE2UR_CALL(zePhysicalMemCreate, (hContext->getZeHandle(), hDevice->ZeDevice, - &PhysMemDesc, &ZePhysMem)); + ze_result_t ZeRes = zePhysicalMemCreate( + hContext->getZeHandle(), hDevice->ZeDevice, &PhysMemDesc, &ZePhysMem); + // Driver has dup'd ImportFdNum internally; close our copy now. + close(ImportFdNum); + + if (ZeRes != ZE_RESULT_SUCCESS) + return ze2urResult(ZeRes); + try { *phPhysMem = new ur_physical_mem_handle_t_(ZePhysMem, hContext, HandleData->Size); diff --git a/unified-runtime/source/adapters/level_zero/physical_mem.hpp b/unified-runtime/source/adapters/level_zero/physical_mem.hpp index 6bb91f0ce268..8f29f47c8c0e 100644 --- a/unified-runtime/source/adapters/level_zero/physical_mem.hpp +++ b/unified-runtime/source/adapters/level_zero/physical_mem.hpp @@ -12,10 +12,13 @@ #include "common/ur_ref_count.hpp" // Opaque handle data exchanged between processes for physical memory IPC. -// Contains the file descriptor exported by zePhysicalMemGetProperties and -// the size of the allocation (required by zePhysicalMemCreate on import). +// Contains the PID and file descriptor exported by zePhysicalMemGetProperties, +// plus the allocation size required by zePhysicalMemCreate on import. +// Cross-process access uses pidfd_getfd(2) (Linux 5.6+): the consumer obtains +// a duplicate of the spawner's DMA-BUF fd via the spawner's PID. struct ZeIPCPhysMemHandleData { - int Fd; + pid_t Pid; // PID of the exporting process + int Fd; // DMA-BUF fd in the exporting process size_t Size; }; diff --git a/unified-runtime/source/adapters/level_zero/v2/memory.cpp b/unified-runtime/source/adapters/level_zero/v2/memory.cpp index 10c789554b3e..77b661cbe459 100644 --- a/unified-runtime/source/adapters/level_zero/v2/memory.cpp +++ b/unified-runtime/source/adapters/level_zero/v2/memory.cpp @@ -9,6 +9,7 @@ #include "memory.hpp" +#include #include #include "../ur_interface_loader.hpp" @@ -960,6 +961,10 @@ ur_result_t urIPCGetPhysMemHandleExp(ur_context_handle_t hContext, return UR_RESULT_ERROR_OUT_OF_HOST_MEMORY; } + // Store the exporting process's PID and fd. The fd stays open until + // urIPCPutPhysMemHandleExp is called. Cross-process consumers use + // pidfd_getfd(2) to obtain their own duplicate of this fd. + HandleData->Pid = getpid(); HandleData->Fd = ExportFd.fd; HandleData->Size = hPhysMem->Size; @@ -1001,11 +1006,31 @@ ur_result_t urIPCOpenPhysMemHandleExp(ur_context_handle_t hContext, auto *HandleData = static_cast(pIPCPhysMemHandleData); + // Obtain a usable fd in the current process. For same-process opens + // (e.g. conformance tests) dup() suffices. For cross-process opens + // use pidfd_getfd(2) which requires the exporting process to be + // ptrace-accessible (e.g. via prctl(PR_SET_PTRACER, PR_SET_PTRACER_ANY)). + int ImportFdNum = -1; + if (HandleData->Pid == getpid()) { + ImportFdNum = dup(HandleData->Fd); + if (ImportFdNum < 0) + return UR_RESULT_ERROR_INVALID_VALUE; + } else { + int PidFd = static_cast(syscall(SYS_pidfd_open, HandleData->Pid, 0)); + if (PidFd < 0) + return UR_RESULT_ERROR_INVALID_VALUE; + ImportFdNum = + static_cast(syscall(SYS_pidfd_getfd, PidFd, HandleData->Fd, 0)); + close(PidFd); + if (ImportFdNum < 0) + return UR_RESULT_ERROR_INVALID_VALUE; + } + ze_external_memory_import_fd_t ImportFd = {}; ImportFd.stype = ZE_STRUCTURE_TYPE_EXTERNAL_MEMORY_IMPORT_FD; ImportFd.pNext = nullptr; ImportFd.flags = ZE_EXTERNAL_MEMORY_TYPE_FLAG_OPAQUE_FD; - ImportFd.fd = HandleData->Fd; + ImportFd.fd = ImportFdNum; ZeStruct PhysMemDesc; PhysMemDesc.pNext = &ImportFd; @@ -1013,8 +1038,14 @@ ur_result_t urIPCOpenPhysMemHandleExp(ur_context_handle_t hContext, PhysMemDesc.size = HandleData->Size; ze_physical_mem_handle_t ZePhysMem; - ZE2UR_CALL(zePhysicalMemCreate, (hContext->getZeHandle(), hDevice->ZeDevice, - &PhysMemDesc, &ZePhysMem)); + ze_result_t ZeRes = zePhysicalMemCreate( + hContext->getZeHandle(), hDevice->ZeDevice, &PhysMemDesc, &ZePhysMem); + // Driver has dup'd ImportFdNum internally; close our copy now. + close(ImportFdNum); + + if (ZeRes != ZE_RESULT_SUCCESS) + return ze2urResult(ZeRes); + try { *phPhysMem = new ur_physical_mem_handle_t_(ZePhysMem, hContext, HandleData->Size); diff --git a/unified-runtime/source/adapters/native_cpu/device.cpp b/unified-runtime/source/adapters/native_cpu/device.cpp index f08de5694a2f..e74668f19728 100644 --- a/unified-runtime/source/adapters/native_cpu/device.cpp +++ b/unified-runtime/source/adapters/native_cpu/device.cpp @@ -381,6 +381,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, static_cast(0)); case UR_DEVICE_INFO_IPC_MEMORY_SUPPORT_EXP: return ReturnValue(false); + case UR_DEVICE_INFO_IPC_PHYSICAL_MEMORY_SUPPORT_EXP: + return ReturnValue(false); case UR_DEVICE_INFO_ATOMIC_FENCE_ORDER_CAPABILITIES: { // Currently for Native CPU fences are implemented using OCK // builtins, so we have different capabilities than atomic operations diff --git a/unified-runtime/source/adapters/opencl/device.cpp b/unified-runtime/source/adapters/opencl/device.cpp index 9654c6b0f764..f8d11ce26c3a 100644 --- a/unified-runtime/source/adapters/opencl/device.cpp +++ b/unified-runtime/source/adapters/opencl/device.cpp @@ -1375,6 +1375,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, } case UR_DEVICE_INFO_IPC_MEMORY_SUPPORT_EXP: return ReturnValue(false); + case UR_DEVICE_INFO_IPC_PHYSICAL_MEMORY_SUPPORT_EXP: + return ReturnValue(false); case UR_DEVICE_INFO_BFLOAT16_CONVERSIONS_NATIVE: { bool Supported = false; UR_RETURN_ON_FAILURE(hDevice->checkDeviceExtensions( From 54bb612fe6ddbed0fd968b63033c778fd4bc58ba Mon Sep 17 00:00:00 2001 From: Lukasz Dorau Date: Fri, 22 May 2026 09:21:50 +0000 Subject: [PATCH 10/17] [UR][L0] Fix refcount bypass in urIPCClosePhysMemHandleExp urIPCClosePhysMemHandleExp was unconditionally calling zePhysicalMemDestroy and deleting the handle object, bypassing the reference count mechanism. If a caller had previously called urPhysicalMemRetain (refcount > 1), the subsequent delete would leave dangling references and a use-after-free on the next urPhysicalMemRelease call. Fix: delegate to ur::level_zero::urPhysicalMemRelease which properly decrements the refcount and only destroys when it reaches zero. Signed-off-by: Lukasz Dorau --- unified-runtime/source/adapters/level_zero/memory.cpp | 8 ++++---- unified-runtime/source/adapters/level_zero/v2/memory.cpp | 8 ++++---- 2 files changed, 8 insertions(+), 8 deletions(-) diff --git a/unified-runtime/source/adapters/level_zero/memory.cpp b/unified-runtime/source/adapters/level_zero/memory.cpp index 473455f08b02..64be9f97969c 100644 --- a/unified-runtime/source/adapters/level_zero/memory.cpp +++ b/unified-runtime/source/adapters/level_zero/memory.cpp @@ -2151,10 +2151,10 @@ ur_result_t urIPCClosePhysMemHandleExp(ur_context_handle_t hContext, if (!hPhysMem) return UR_RESULT_ERROR_INVALID_NULL_HANDLE; - ur_result_t Res = ze2urResult( - zePhysicalMemDestroy(hContext->getZeHandle(), hPhysMem->ZePhysicalMem)); - delete hPhysMem; - return Res; + // Delegate to urPhysicalMemRelease so the refcount is respected: if the + // handle has been retained (refcount > 1) it will not be destroyed until + // all references are released. + return ur::level_zero::urPhysicalMemRelease(hPhysMem); } } // namespace ur::level_zero diff --git a/unified-runtime/source/adapters/level_zero/v2/memory.cpp b/unified-runtime/source/adapters/level_zero/v2/memory.cpp index 77b661cbe459..1ea69754aecc 100644 --- a/unified-runtime/source/adapters/level_zero/v2/memory.cpp +++ b/unified-runtime/source/adapters/level_zero/v2/memory.cpp @@ -1066,10 +1066,10 @@ ur_result_t urIPCClosePhysMemHandleExp(ur_context_handle_t hContext, if (!hPhysMem) return UR_RESULT_ERROR_INVALID_NULL_HANDLE; - ur_result_t Res = ze2urResult( - zePhysicalMemDestroy(hContext->getZeHandle(), hPhysMem->ZePhysicalMem)); - delete hPhysMem; - return Res; + // Delegate to urPhysicalMemRelease so the refcount is respected: if the + // handle has been retained (refcount > 1) it will not be destroyed until + // all references are released. + return ur::level_zero::urPhysicalMemRelease(hPhysMem); } } // namespace ur::level_zero From 8c444fe5ca5c50bf4e033f9e44fd6dc9f44dc587 Mon Sep 17 00:00:00 2001 From: Lukasz Dorau Date: Fri, 22 May 2026 09:52:09 +0000 Subject: [PATCH 11/17] [UR][test] Fix wrong device capability check in IPC physical mem fixture The urIPCPhysMemTest fixture was querying UR_DEVICE_INFO_IPC_MEMORY_SUPPORT_EXP (USM allocation IPC) instead of UR_DEVICE_INFO_IPC_PHYSICAL_MEMORY_SUPPORT_EXP to guard the IPC physical_mem tests. These are distinct capabilities and could diverge on future adapters, causing tests to run on unsupported devices or to be incorrectly skipped. Also update the skip message to be specific about physical memory. Signed-off-by: Lukasz Dorau --- .../virtual_memory/urIPCPhysMemHandleExpFixtures.hpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/unified-runtime/test/conformance/virtual_memory/urIPCPhysMemHandleExpFixtures.hpp b/unified-runtime/test/conformance/virtual_memory/urIPCPhysMemHandleExpFixtures.hpp index 4134ec097d49..93430e0dacc8 100644 --- a/unified-runtime/test/conformance/virtual_memory/urIPCPhysMemHandleExpFixtures.hpp +++ b/unified-runtime/test/conformance/virtual_memory/urIPCPhysMemHandleExpFixtures.hpp @@ -14,11 +14,11 @@ struct urIPCPhysMemTest : uur::urVirtualMemGranularityTest { size = granularity * 256; ur_bool_t ipc_support = false; - ASSERT_SUCCESS(urDeviceGetInfo(device, - UR_DEVICE_INFO_IPC_MEMORY_SUPPORT_EXP, - sizeof(ur_bool_t), &ipc_support, nullptr)); + ASSERT_SUCCESS( + urDeviceGetInfo(device, UR_DEVICE_INFO_IPC_PHYSICAL_MEMORY_SUPPORT_EXP, + sizeof(ur_bool_t), &ipc_support, nullptr)); if (!ipc_support) { - GTEST_SKIP() << "IPC memory is not supported."; + GTEST_SKIP() << "IPC physical memory is not supported."; } ur_physical_mem_properties_t properties{ From b9e281a0503da2ec6b7fdb8e0f7a5dccb9d965f0 Mon Sep 17 00:00:00 2001 From: Lukasz Dorau Date: Fri, 22 May 2026 11:35:42 +0000 Subject: [PATCH 12/17] [UR][L0][SYCL] Fix urPhysicalMemGetInfo and IPC-opened physical_mem size Two related fixes: 1. urPhysicalMemGetInfo (L0 adapter) only returned REFERENCE_COUNT; all other queries (CONTEXT, DEVICE, SIZE, PROPERTIES) returned UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION despite the data being available in the handle struct. Fix: add Device and EnableIpc fields to ur_physical_mem_handle_t_, populate them in urPhysicalMemCreate and urIPCOpenPhysMemHandleExp, and implement all five info cases in urPhysicalMemGetInfo. 2. physical_mem objects created via urIPCOpenPhysMemHandleExp were constructed with NumBytes=0, making physical_mem::size() return 0 instead of the actual allocation size. Fix: after opening, query UR_PHYSICAL_MEM_INFO_SIZE and pass the result to the physical_mem_impl constructor. Signed-off-by: Lukasz Dorau --- sycl/source/ipc_memory.cpp | 9 +++++++- .../source/adapters/level_zero/memory.cpp | 5 +++-- .../adapters/level_zero/physical_mem.cpp | 21 +++++++++++++++---- .../adapters/level_zero/physical_mem.hpp | 13 ++++++++++-- .../source/adapters/level_zero/v2/memory.cpp | 5 +++-- 5 files changed, 42 insertions(+), 11 deletions(-) diff --git a/sycl/source/ipc_memory.cpp b/sycl/source/ipc_memory.cpp index 89af391164af..fb83aecbf121 100644 --- a/sycl/source/ipc_memory.cpp +++ b/sycl/source/ipc_memory.cpp @@ -88,8 +88,15 @@ openIPCPhysMemHandle(const std::byte *HandleData, size_t HandleDataSize, "urIPCOpenPhysMemHandleExp returned success but did not produce a " "valid physical memory handle."); + // 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=*/0, PhysMemHandle); + *getSyclObjImpl(Dev), Ctx, NumBytes, PhysMemHandle); return sycl::detail::createSyclObjFromImpl< ext::oneapi::experimental::physical_mem>(PhysMemImpl); } diff --git a/unified-runtime/source/adapters/level_zero/memory.cpp b/unified-runtime/source/adapters/level_zero/memory.cpp index 64be9f97969c..0a8b74cbfe33 100644 --- a/unified-runtime/source/adapters/level_zero/memory.cpp +++ b/unified-runtime/source/adapters/level_zero/memory.cpp @@ -2132,8 +2132,9 @@ ur_result_t urIPCOpenPhysMemHandleExp(ur_context_handle_t hContext, return ze2urResult(ZeRes); try { - *phPhysMem = - new ur_physical_mem_handle_t_(ZePhysMem, hContext, HandleData->Size); + *phPhysMem = new ur_physical_mem_handle_t_(ZePhysMem, hContext, hDevice, + HandleData->Size, + /*EnableIpc=*/false); } catch (const std::bad_alloc &) { zePhysicalMemDestroy(hContext->getZeHandle(), ZePhysMem); return UR_RESULT_ERROR_OUT_OF_HOST_MEMORY; diff --git a/unified-runtime/source/adapters/level_zero/physical_mem.cpp b/unified-runtime/source/adapters/level_zero/physical_mem.cpp index 8f3010d1b683..563476415825 100644 --- a/unified-runtime/source/adapters/level_zero/physical_mem.cpp +++ b/unified-runtime/source/adapters/level_zero/physical_mem.cpp @@ -42,8 +42,8 @@ ur_result_t urPhysicalMemCreate( ZE2UR_CALL(zePhysicalMemCreate, (hContext->getZeHandle(), hDevice->ZeDevice, &PhysicalMemDesc, &ZePhysicalMem)); try { - *phPhysicalMem = - new ur_physical_mem_handle_t_(ZePhysicalMem, hContext, size); + *phPhysicalMem = new ur_physical_mem_handle_t_(ZePhysicalMem, hContext, + hDevice, size, EnableIpc); } catch (const std::bad_alloc &) { zePhysicalMemDestroy(hContext->getZeHandle(), ZePhysicalMem); return UR_RESULT_ERROR_OUT_OF_HOST_MEMORY; @@ -80,9 +80,22 @@ ur_result_t urPhysicalMemGetInfo(ur_physical_mem_handle_t hPhysicalMem, UrReturnHelper ReturnValue(propSize, pPropValue, pPropSizeRet); switch (propName) { - case UR_PHYSICAL_MEM_INFO_REFERENCE_COUNT: { - return ReturnValue(hPhysicalMem->RefCount.getCount()); + case UR_PHYSICAL_MEM_INFO_CONTEXT: + return ReturnValue(hPhysicalMem->Context); + case UR_PHYSICAL_MEM_INFO_DEVICE: + return ReturnValue(hPhysicalMem->Device); + case UR_PHYSICAL_MEM_INFO_SIZE: + return ReturnValue(hPhysicalMem->Size); + case UR_PHYSICAL_MEM_INFO_PROPERTIES: { + ur_physical_mem_flags_t Flags = static_cast(0); + if (hPhysicalMem->EnableIpc) + Flags = UR_PHYSICAL_MEM_FLAG_ENABLE_IPC; + ur_physical_mem_properties_t Props = { + UR_STRUCTURE_TYPE_PHYSICAL_MEM_PROPERTIES, nullptr, Flags}; + return ReturnValue(Props); } + case UR_PHYSICAL_MEM_INFO_REFERENCE_COUNT: + return ReturnValue(hPhysicalMem->RefCount.getCount()); default: return UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION; } diff --git a/unified-runtime/source/adapters/level_zero/physical_mem.hpp b/unified-runtime/source/adapters/level_zero/physical_mem.hpp index 8f29f47c8c0e..99db18a83a39 100644 --- a/unified-runtime/source/adapters/level_zero/physical_mem.hpp +++ b/unified-runtime/source/adapters/level_zero/physical_mem.hpp @@ -24,8 +24,11 @@ struct ZeIPCPhysMemHandleData { struct ur_physical_mem_handle_t_ : ur_object { ur_physical_mem_handle_t_(ze_physical_mem_handle_t ZePhysicalMem, - ur_context_handle_t Context, size_t Size) - : ZePhysicalMem{ZePhysicalMem}, Context{Context}, Size{Size} {} + ur_context_handle_t Context, + ur_device_handle_t Device, size_t Size, + bool EnableIpc) + : ZePhysicalMem{ZePhysicalMem}, Context{Context}, Device{Device}, + Size{Size}, EnableIpc{EnableIpc} {} // Level Zero physical memory handle. ze_physical_mem_handle_t ZePhysicalMem; @@ -33,8 +36,14 @@ struct ur_physical_mem_handle_t_ : ur_object { // Keeps the PI context of this memory handle. ur_context_handle_t Context; + // Device this physical memory was allocated on. + ur_device_handle_t Device; + // Size in bytes of this physical memory allocation. size_t Size; + // Whether this allocation was created with IPC export enabled. + bool EnableIpc; + ur::RefCount RefCount; }; diff --git a/unified-runtime/source/adapters/level_zero/v2/memory.cpp b/unified-runtime/source/adapters/level_zero/v2/memory.cpp index 1ea69754aecc..346432fbba54 100644 --- a/unified-runtime/source/adapters/level_zero/v2/memory.cpp +++ b/unified-runtime/source/adapters/level_zero/v2/memory.cpp @@ -1047,8 +1047,9 @@ ur_result_t urIPCOpenPhysMemHandleExp(ur_context_handle_t hContext, return ze2urResult(ZeRes); try { - *phPhysMem = - new ur_physical_mem_handle_t_(ZePhysMem, hContext, HandleData->Size); + *phPhysMem = new ur_physical_mem_handle_t_(ZePhysMem, hContext, hDevice, + HandleData->Size, + /*EnableIpc=*/false); } catch (const std::bad_alloc &) { zePhysicalMemDestroy(hContext->getZeHandle(), ZePhysMem); return UR_RESULT_ERROR_OUT_OF_HOST_MEMORY; From 0df780be289f4ce04c3ad4de55bb6189ba695424 Mon Sep 17 00:00:00 2001 From: Lukasz Dorau Date: Fri, 22 May 2026 13:18:05 +0000 Subject: [PATCH 13/17] [SYCL] Fix handle leak in openIPCPhysMemHandle on exception After urIPCOpenPhysMemHandleExp succeeds, a failure in the subsequent urPhysicalMemGetInfo call (or std::bad_alloc in make_shared) would throw without releasing PhysMemHandle, leaking the GPU resource. Wrap the post-open work in a try-catch that calls urIPCClosePhysMemHandleExp on any exception before re-throwing, ensuring the handle is always released. Signed-off-by: Lukasz Dorau --- sycl/source/ipc_memory.cpp | 31 ++++++++++++++++++++----------- 1 file changed, 20 insertions(+), 11 deletions(-) diff --git a/sycl/source/ipc_memory.cpp b/sycl/source/ipc_memory.cpp index fb83aecbf121..f59b32152777 100644 --- a/sycl/source/ipc_memory.cpp +++ b/sycl/source/ipc_memory.cpp @@ -88,17 +88,26 @@ openIPCPhysMemHandle(const std::byte *HandleData, size_t HandleDataSize, "urIPCOpenPhysMemHandleExp returned success but did not produce a " "valid physical memory handle."); - // 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< - ext::oneapi::experimental::physical_mem>(PhysMemImpl); + // Any failure after this point must release PhysMemHandle to avoid leaking + // the GPU resource. Wrap in try-catch so both urPhysicalMemGetInfo failures + // and std::bad_alloc from make_shared are handled. + 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< + ext::oneapi::experimental::physical_mem>(PhysMemImpl); + } catch (...) { + Adapter.call_nocheck( + CtxImpl->getHandleRef(), PhysMemHandle); + throw; + } } } // namespace detail From fb587b350bcef7d8e992183eb74f45c123136226 Mon Sep 17 00:00:00 2001 From: Lukasz Dorau Date: Fri, 22 May 2026 15:01:51 +0000 Subject: [PATCH 14/17] [UR][SYCL] Fix CI build failures for IPC physical_mem - Move UR_DEVICE_INFO_IPC_PHYSICAL_MEMORY_SUPPORT_EXP enum extension from virtual_memory.yml to exp-inter-process-communication.yml where all other IPC-related spec content lives - Regenerate tools/urinfo/urinfo.hpp to include the new device info entry (fixes 'check-generated' CI failure on Windows) - Replace mixed-type ternary with if/else in physical_mem_impl.hpp to fix -Werror=extra: 'enumerated and non-enumerated type in conditional expression' (fixes SYCL build CI failure) --- sycl/source/detail/physical_mem_impl.hpp | 7 ++++--- .../scripts/core/exp-inter-process-communication.yml | 10 ++++++++++ unified-runtime/scripts/core/virtual_memory.yml | 11 ----------- unified-runtime/tools/urinfo/urinfo.hpp | 3 +++ 4 files changed, 17 insertions(+), 14 deletions(-) diff --git a/sycl/source/detail/physical_mem_impl.hpp b/sycl/source/detail/physical_mem_impl.hpp index e330a13e8faa..c71ab06f200a 100644 --- a/sycl/source/detail/physical_mem_impl.hpp +++ b/sycl/source/detail/physical_mem_impl.hpp @@ -43,10 +43,11 @@ class physical_mem_impl { MNumBytes(NumBytes), MEnabledIpc(EnableIpc) { adapter_impl &Adapter = MContext->getAdapter(); + ur_physical_mem_flags_t Flags = ur_physical_mem_flags_t(0); + if (EnableIpc) + Flags = UR_PHYSICAL_MEM_FLAG_ENABLE_IPC; ur_physical_mem_properties_t Props = { - UR_STRUCTURE_TYPE_PHYSICAL_MEM_PROPERTIES, nullptr, - EnableIpc ? UR_PHYSICAL_MEM_FLAG_ENABLE_IPC - : ur_physical_mem_flags_t(0)}; + UR_STRUCTURE_TYPE_PHYSICAL_MEM_PROPERTIES, nullptr, Flags}; auto Err = Adapter.call_nocheck( MContext->getHandleRef(), MDevice.getHandleRef(), MNumBytes, &Props, diff --git a/unified-runtime/scripts/core/exp-inter-process-communication.yml b/unified-runtime/scripts/core/exp-inter-process-communication.yml index 312dba7919ad..0c2eacf99b6d 100644 --- a/unified-runtime/scripts/core/exp-inter-process-communication.yml +++ b/unified-runtime/scripts/core/exp-inter-process-communication.yml @@ -23,6 +23,16 @@ etors: value: "0x2023" desc: "[$x_bool_t] returns true if the device supports inter-process communicable memory handles" --- #-------------------------------------------------------------------------- +type: enum +extend: true +typed_etors: true +desc: "Extension enums to $x_device_info_t to support inter-process communicable physical memory handles." +name: $x_device_info_t +etors: + - name: IPC_PHYSICAL_MEMORY_SUPPORT_EXP + value: "0x2024" + desc: "[$x_bool_t] returns true if the device supports inter-process communicable physical memory handles" +--- #-------------------------------------------------------------------------- type: function desc: "Gets an inter-process memory handle for a pointer to device USM memory" class: $xIPC diff --git a/unified-runtime/scripts/core/virtual_memory.yml b/unified-runtime/scripts/core/virtual_memory.yml index d1b841099486..4fe51b74dd55 100644 --- a/unified-runtime/scripts/core/virtual_memory.yml +++ b/unified-runtime/scripts/core/virtual_memory.yml @@ -11,17 +11,6 @@ type: header desc: "Intel $OneApi Unified Runtime APIs" ordinal: "4" ---- #-------------------------------------------------------------------------- -type: enum -extend: true -typed_etors: true -desc: "Extension enums to $x_device_info_t to support inter-process communicable physical memory handles." -name: $x_device_info_t -etors: - - name: IPC_PHYSICAL_MEMORY_SUPPORT_EXP - value: "0x2024" - desc: "[$x_bool_t] returns true if the device supports inter-process communicable physical memory handles" - --- #-------------------------------------------------------------------------- type: enum desc: "Virtual memory granularity info" diff --git a/unified-runtime/tools/urinfo/urinfo.hpp b/unified-runtime/tools/urinfo/urinfo.hpp index d5b15e5aaa00..fe99cf3ef5f0 100644 --- a/unified-runtime/tools/urinfo/urinfo.hpp +++ b/unified-runtime/tools/urinfo/urinfo.hpp @@ -466,6 +466,9 @@ inline void printDeviceInfos(ur_device_handle_t hDevice, std::cout << prefix; printDeviceInfo(hDevice, UR_DEVICE_INFO_IPC_MEMORY_SUPPORT_EXP); std::cout << prefix; + printDeviceInfo(hDevice, + UR_DEVICE_INFO_IPC_PHYSICAL_MEMORY_SUPPORT_EXP); + std::cout << prefix; printDeviceInfo(hDevice, UR_DEVICE_INFO_ASYNC_USM_ALLOCATIONS_SUPPORT_EXP); std::cout << prefix; From 6f41d314b1160a8ab5d0d595e76a5e6291eb0abb Mon Sep 17 00:00:00 2001 From: Lukasz Dorau Date: Fri, 22 May 2026 15:27:17 +0000 Subject: [PATCH 15/17] [UR][OpenCL] Fix clang-format for opencl/common.hpp Restore clang-format 20 correct indentation for CL_API_ENTRY function pointer type declarations that was accidentally reverted during rebase. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- .../source/adapters/opencl/common.hpp | 69 ++++++++++--------- 1 file changed, 35 insertions(+), 34 deletions(-) diff --git a/unified-runtime/source/adapters/opencl/common.hpp b/unified-runtime/source/adapters/opencl/common.hpp index fadab9d19992..6644b42121cc 100644 --- a/unified-runtime/source/adapters/opencl/common.hpp +++ b/unified-runtime/source/adapters/opencl/common.hpp @@ -224,37 +224,37 @@ CONSTFIX char GetKernelSubGroupInfoName[] = "clGetKernelSubGroupInfoKHR"; #undef CONSTFIX using clGetDeviceFunctionPointerINTEL_fn = CL_API_ENTRY -cl_int(CL_API_CALL *)(cl_device_id device, cl_program program, - const char *FuncName, cl_ulong *ret_ptr); + cl_int(CL_API_CALL *)(cl_device_id device, cl_program program, + const char *FuncName, cl_ulong *ret_ptr); using clGetDeviceGlobalVariablePointerINTEL_fn = CL_API_ENTRY -cl_int(CL_API_CALL *)(cl_device_id device, cl_program program, - const char *globalVariableName, - size_t *globalVariableSizeRet, - void **globalVariablePointerRet); + cl_int(CL_API_CALL *)(cl_device_id device, cl_program program, + const char *globalVariableName, + size_t *globalVariableSizeRet, + void **globalVariablePointerRet); using clEnqueueWriteGlobalVariableINTEL_fn = CL_API_ENTRY -cl_int(CL_API_CALL *)(cl_command_queue, cl_program, const char *, cl_bool, - size_t, size_t, const void *, cl_uint, const cl_event *, - cl_event *); + cl_int(CL_API_CALL *)(cl_command_queue, cl_program, const char *, cl_bool, + size_t, size_t, const void *, cl_uint, + const cl_event *, cl_event *); using clEnqueueReadGlobalVariableINTEL_fn = CL_API_ENTRY -cl_int(CL_API_CALL *)(cl_command_queue, cl_program, const char *, cl_bool, - size_t, size_t, void *, cl_uint, const cl_event *, - cl_event *); + cl_int(CL_API_CALL *)(cl_command_queue, cl_program, const char *, cl_bool, + size_t, size_t, void *, cl_uint, const cl_event *, + cl_event *); using clEnqueueReadHostPipeINTEL_fn = CL_API_ENTRY -cl_int(CL_API_CALL *)(cl_command_queue queue, cl_program program, - const char *pipe_symbol, cl_bool blocking, void *ptr, - size_t size, cl_uint num_events_in_waitlist, - const cl_event *events_waitlist, cl_event *event); + cl_int(CL_API_CALL *)(cl_command_queue queue, cl_program program, + const char *pipe_symbol, cl_bool blocking, void *ptr, + size_t size, cl_uint num_events_in_waitlist, + const cl_event *events_waitlist, cl_event *event); using clEnqueueWriteHostPipeINTEL_fn = CL_API_ENTRY -cl_int(CL_API_CALL *)(cl_command_queue queue, cl_program program, - const char *pipe_symbol, cl_bool blocking, - const void *ptr, size_t size, - cl_uint num_events_in_waitlist, - const cl_event *events_waitlist, cl_event *event); + cl_int(CL_API_CALL *)(cl_command_queue queue, cl_program program, + const char *pipe_symbol, cl_bool blocking, + const void *ptr, size_t size, + cl_uint num_events_in_waitlist, + const cl_event *events_waitlist, cl_event *event); using clCreateCommandBufferKHR_fn = CL_API_ENTRY cl_command_buffer_khr( CL_API_CALL *)(cl_uint num_queues, const cl_command_queue *queues, @@ -262,13 +262,13 @@ using clCreateCommandBufferKHR_fn = CL_API_ENTRY cl_command_buffer_khr( cl_int *errcode_ret); using clRetainCommandBufferKHR_fn = CL_API_ENTRY -cl_int(CL_API_CALL *)(cl_command_buffer_khr command_buffer); + cl_int(CL_API_CALL *)(cl_command_buffer_khr command_buffer); using clReleaseCommandBufferKHR_fn = CL_API_ENTRY -cl_int(CL_API_CALL *)(cl_command_buffer_khr command_buffer); + cl_int(CL_API_CALL *)(cl_command_buffer_khr command_buffer); using clFinalizeCommandBufferKHR_fn = CL_API_ENTRY -cl_int(CL_API_CALL *)(cl_command_buffer_khr command_buffer); + cl_int(CL_API_CALL *)(cl_command_buffer_khr command_buffer); using clCommandNDRangeKernelKHR_fn = CL_API_ENTRY cl_int(CL_API_CALL *)( cl_command_buffer_khr command_buffer, cl_command_queue command_queue, @@ -328,26 +328,27 @@ using clCommandSVMMemFillKHR_fn = CL_API_ENTRY cl_int(CL_API_CALL *)( cl_sync_point_khr *sync_point, cl_mutable_command_khr *mutable_handle); using clEnqueueCommandBufferKHR_fn = CL_API_ENTRY -cl_int(CL_API_CALL *)(cl_uint num_queues, cl_command_queue *queues, - cl_command_buffer_khr command_buffer, - cl_uint num_events_in_wait_list, - const cl_event *event_wait_list, cl_event *event); + cl_int(CL_API_CALL *)(cl_uint num_queues, cl_command_queue *queues, + cl_command_buffer_khr command_buffer, + cl_uint num_events_in_wait_list, + const cl_event *event_wait_list, cl_event *event); using clGetCommandBufferInfoKHR_fn = CL_API_ENTRY cl_int(CL_API_CALL *)( cl_command_buffer_khr command_buffer, cl_command_buffer_info_khr param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret); using clUpdateMutableCommandsKHR_fn = CL_API_ENTRY -cl_int(CL_API_CALL *)(cl_command_buffer_khr command_buffer, cl_uint num_configs, - const cl_command_buffer_update_type_khr *config_types, - const void **configs); + cl_int(CL_API_CALL *)(cl_command_buffer_khr command_buffer, + cl_uint num_configs, + const cl_command_buffer_update_type_khr *config_types, + const void **configs); using clCreateProgramWithILKHR_fn = CL_API_ENTRY -cl_program(CL_API_CALL *)(cl_context, const void *, size_t, cl_int *); + cl_program(CL_API_CALL *)(cl_context, const void *, size_t, cl_int *); using clGetKernelSubGroupInfoKHR_fn = CL_API_ENTRY -cl_int(CL_API_CALL *)(cl_kernel, cl_device_id, cl_kernel_sub_group_info, size_t, - const void *, size_t, void *, size_t *); + cl_int(CL_API_CALL *)(cl_kernel, cl_device_id, cl_kernel_sub_group_info, + size_t, const void *, size_t, void *, size_t *); template struct FuncPtrCache { std::map Map; From b23272afa968ee5e49a3283e1f6824a7a3c37021 Mon Sep 17 00:00:00 2001 From: Lukasz Dorau Date: Fri, 22 May 2026 15:54:42 +0000 Subject: [PATCH 16/17] [UR][SYCL][L0] Fix CI failures for IPC physical_mem - Guard ZeIPCPhysMemHandleData struct and Linux-specific IPC function bodies in L0 v1 and v2 adapters with #ifdef __linux__ to fix Windows build errors (pid_t and sys/syscall.h are Linux-only). - Add ext_oneapi_ipc_physical_memory aspect to DeviceConfigFile.td to fix device_config_file_aspects LIT test failure. - Update sycl_symbols_linux.dump ABI reference: replace old 3-param physical_mem constructor with new 4-param one (adds bool EnableIpc), and add all new IPC physical_memory API symbols. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- .../llvm/SYCLLowerIR/DeviceConfigFile.td | 2 ++ sycl/test/abi/sycl_symbols_linux.dump | 19 ++++++++------ .../source/adapters/level_zero/memory.cpp | 25 +++++++++++++++++++ .../adapters/level_zero/physical_mem.hpp | 3 +++ .../source/adapters/level_zero/v2/memory.cpp | 25 +++++++++++++++++++ 5 files changed, 66 insertions(+), 8 deletions(-) diff --git a/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td b/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td index 262a9dce6535..057cd0648856 100644 --- a/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td +++ b/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td @@ -97,6 +97,7 @@ def Aspectext_oneapi_clock_work_group : Aspect<"ext_oneapi_clock_work_group">; def Aspectext_oneapi_clock_device : Aspect<"ext_oneapi_clock_device">; def Aspectext_oneapi_is_integrated_gpu : Aspect<"ext_oneapi_is_integrated_gpu">; def Aspectext_oneapi_ipc_memory : Aspect<"ext_oneapi_ipc_memory">; +def Aspectext_oneapi_ipc_physical_memory : Aspect<"ext_oneapi_ipc_physical_memory">; def Aspectext_oneapi_device_wait : Aspect<"ext_oneapi_device_wait">; def Aspectext_intel_xe_stack_count : Aspect<"ext_intel_xe_stack_count">; def Aspectext_intel_xe_regions_per_stack : Aspect<"ext_intel_xe_regions_per_stack">; @@ -183,6 +184,7 @@ def : TargetInfo<"__TestAspectList", Aspectext_oneapi_clock_device, Aspectext_oneapi_is_integrated_gpu, Aspectext_oneapi_ipc_memory, + Aspectext_oneapi_ipc_physical_memory, Aspectext_oneapi_device_wait, Aspectext_intel_xe_stack_count, Aspectext_intel_xe_regions_per_stack, diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 02ff46c47878..d93a8eea57c5 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -1,12 +1,10 @@ -################################################################################ -# This file is automatically generated by abi_check.py tool. -# DO NOT EDIT IT MANUALLY. Refer to sycl/doc/developer/ABIPolicyGuide.md for more info. -################################################################################ -# RUN: env LLVM_BIN_PATH=%llvm_build_bin_dir %python %sycl_tools_src_dir/abi_check.py --mode check_symbols --reference %s %sycl_libs_dir/libsycl.so +# DO NOT EDIT IT MANUALLY. Refer to sycl/doc/developer/ABIPolicyGuide.md for more info. # REQUIRES: linux +# RUN: env LLVM_BIN_PATH=%llvm_build_bin_dir %python %sycl_tools_src_dir/abi_check.py --mode check_symbols --reference %s %sycl_libs_dir/libsycl.so +# This file is automatically generated by abi_check.py tool. # UNSUPPORTED: libcxx - +################################################################################ _ZN4sycl3_V110__abs_implENS0_3vecIaLi16EEE _ZN4sycl3_V110__abs_implENS0_3vecIaLi1EEE _ZN4sycl3_V110__abs_implENS0_3vecIaLi2EEE @@ -3006,8 +3004,8 @@ _ZN4sycl3_V13ext6oneapi12experimental12create_imageERNS3_9image_memERKNS3_16imag _ZN4sycl3_V13ext6oneapi12experimental12create_imageERNS3_9image_memERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental12create_imageERNS3_9image_memERKNS3_22bindless_image_samplerERKNS3_16image_descriptorERKNS0_5queueE _ZN4sycl3_V13ext6oneapi12experimental12create_imageERNS3_9image_memERKNS3_22bindless_image_samplerERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE -_ZN4sycl3_V13ext6oneapi12experimental12physical_memC1ERKNS0_6deviceERKNS0_7contextEm -_ZN4sycl3_V13ext6oneapi12experimental12physical_memC2ERKNS0_6deviceERKNS0_7contextEm +_ZN4sycl3_V13ext6oneapi12experimental12physical_memC1ERKNS0_6deviceERKNS0_7contextEmb +_ZN4sycl3_V13ext6oneapi12experimental12physical_memC2ERKNS0_6deviceERKNS0_7contextEmb _ZN4sycl3_V13ext6oneapi12experimental13aligned_allocEmmRKNS0_6deviceENS0_3usm5allocERKNS0_13property_listE _ZN4sycl3_V13ext6oneapi12experimental13malloc_deviceEmRKNS0_6deviceERKNS0_13property_listE _ZN4sycl3_V13ext6oneapi12experimental13malloc_sharedEmRKNS0_6deviceERKNS0_13property_listE @@ -3076,6 +3074,9 @@ _ZN4sycl3_V13ext6oneapi12experimental27alloc_exportable_device_memEmmNS3_24exter _ZN4sycl3_V13ext6oneapi12experimental27unmap_external_image_memoryENS3_16image_mem_handleENS3_10image_typeERKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental28unmap_external_linear_memoryEPvRKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental30supports_importing_handle_typeENS3_24external_mem_handle_typeERKNS0_6deviceE +_ZN4sycl3_V13ext6oneapi12experimental3ipc15physical_memory3getERKNS3_12physical_memE +_ZN4sycl3_V13ext6oneapi12experimental3ipc15physical_memory3putERNS4_6handleERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental3ipc15physical_memory4openERKSt6vectorISt4byteSaIS7_EERKNS0_7contextERKNS0_6deviceE _ZN4sycl3_V13ext6oneapi12experimental3ipc6memory3getEPvRKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental3ipc6memory3putERNS4_6handleERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental3ipc6memory5closeEPvRKNS0_7contextE @@ -3319,6 +3320,7 @@ _ZN4sycl3_V16detail20get_kernel_info_implINS0_4info22kernel_device_specific22com _ZN4sycl3_V16detail20get_kernel_info_implINS0_4info22kernel_device_specific23compile_work_group_sizeEEENT_11return_typeERNS1_12context_implERNS1_11device_implERNS1_16DeviceKernelInfoE _ZN4sycl3_V16detail20get_kernel_info_implINS0_4info22kernel_device_specific34preferred_work_group_size_multipleEEENT_11return_typeERNS1_12context_implERNS1_11device_implERNS1_16DeviceKernelInfoE _ZN4sycl3_V16detail20markBufferAsInternalERKSt10shared_ptrINS1_11buffer_implEE +_ZN4sycl3_V16detail20openIPCPhysMemHandleEPKSt4bytemRKNS0_7contextERKNS0_6deviceE _ZN4sycl3_V16detail20verifyReductionPropsERKNS0_13property_listE _ZN4sycl3_V16detail21LocalAccessorBaseHost12getNumOfDimsEv _ZN4sycl3_V16detail21LocalAccessorBaseHost14getElementSizeEv @@ -3648,6 +3650,7 @@ _ZNK4sycl3_V13ext6oneapi12experimental11memory_pool21get_used_size_currentEv _ZNK4sycl3_V13ext6oneapi12experimental11memory_pool25get_reserved_size_currentEv _ZNK4sycl3_V13ext6oneapi12experimental12physical_mem10get_deviceEv _ZNK4sycl3_V13ext6oneapi12experimental12physical_mem11get_contextEv +_ZNK4sycl3_V13ext6oneapi12experimental12physical_mem22ext_oneapi_ipc_enabledEv _ZNK4sycl3_V13ext6oneapi12experimental12physical_mem3mapEmmNS3_19address_access_modeEm _ZNK4sycl3_V13ext6oneapi12experimental12physical_mem4sizeEv _ZNK4sycl3_V13ext6oneapi12experimental21dynamic_command_group16get_active_indexEv diff --git a/unified-runtime/source/adapters/level_zero/memory.cpp b/unified-runtime/source/adapters/level_zero/memory.cpp index 0a8b74cbfe33..bec36b889240 100644 --- a/unified-runtime/source/adapters/level_zero/memory.cpp +++ b/unified-runtime/source/adapters/level_zero/memory.cpp @@ -10,8 +10,10 @@ #include #include #include +#ifdef __linux__ #include #include +#endif #include #include "context.hpp" @@ -2018,6 +2020,7 @@ ur_result_t urIPCGetPhysMemHandleExp(ur_context_handle_t hContext, ur_physical_mem_handle_t hPhysMem, void **ppIPCPhysMemHandleData, size_t *pIPCPhysMemHandleDataSizeRet) { +#ifdef __linux__ if (!hContext) return UR_RESULT_ERROR_INVALID_NULL_HANDLE; if (!hPhysMem) @@ -2056,10 +2059,18 @@ ur_result_t urIPCGetPhysMemHandleExp(ur_context_handle_t hContext, *ppIPCPhysMemHandleData = HandleData; *pIPCPhysMemHandleDataSizeRet = sizeof(ZeIPCPhysMemHandleData); return UR_RESULT_SUCCESS; +#else + (void)hContext; + (void)hPhysMem; + (void)ppIPCPhysMemHandleData; + (void)pIPCPhysMemHandleDataSizeRet; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +#endif // __linux__ } ur_result_t urIPCPutPhysMemHandleExp(ur_context_handle_t hContext, void *pIPCPhysMemHandleData) { +#ifdef __linux__ if (!hContext) return UR_RESULT_ERROR_INVALID_NULL_HANDLE; if (!pIPCPhysMemHandleData) @@ -2070,6 +2081,11 @@ ur_result_t urIPCPutPhysMemHandleExp(ur_context_handle_t hContext, close(HandleData->Fd); delete HandleData; return UR_RESULT_SUCCESS; +#else + (void)hContext; + (void)pIPCPhysMemHandleData; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +#endif // __linux__ } ur_result_t urIPCOpenPhysMemHandleExp(ur_context_handle_t hContext, @@ -2077,6 +2093,7 @@ ur_result_t urIPCOpenPhysMemHandleExp(ur_context_handle_t hContext, void *pIPCPhysMemHandleData, size_t ipcPhysMemHandleDataSize, ur_physical_mem_handle_t *phPhysMem) { +#ifdef __linux__ if (!hContext) return UR_RESULT_ERROR_INVALID_NULL_HANDLE; if (!hDevice) @@ -2143,6 +2160,14 @@ ur_result_t urIPCOpenPhysMemHandleExp(ur_context_handle_t hContext, return UR_RESULT_ERROR_UNKNOWN; } return UR_RESULT_SUCCESS; +#else + (void)hContext; + (void)hDevice; + (void)pIPCPhysMemHandleData; + (void)ipcPhysMemHandleDataSize; + (void)phPhysMem; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +#endif // __linux__ } ur_result_t urIPCClosePhysMemHandleExp(ur_context_handle_t hContext, diff --git a/unified-runtime/source/adapters/level_zero/physical_mem.hpp b/unified-runtime/source/adapters/level_zero/physical_mem.hpp index 99db18a83a39..290d6c826f46 100644 --- a/unified-runtime/source/adapters/level_zero/physical_mem.hpp +++ b/unified-runtime/source/adapters/level_zero/physical_mem.hpp @@ -16,11 +16,14 @@ // plus the allocation size required by zePhysicalMemCreate on import. // Cross-process access uses pidfd_getfd(2) (Linux 5.6+): the consumer obtains // a duplicate of the spawner's DMA-BUF fd via the spawner's PID. +// Only defined on Linux because pid_t and DMA-BUF fds are Linux concepts. +#ifdef __linux__ struct ZeIPCPhysMemHandleData { pid_t Pid; // PID of the exporting process int Fd; // DMA-BUF fd in the exporting process size_t Size; }; +#endif // __linux__ struct ur_physical_mem_handle_t_ : ur_object { ur_physical_mem_handle_t_(ze_physical_mem_handle_t ZePhysicalMem, diff --git a/unified-runtime/source/adapters/level_zero/v2/memory.cpp b/unified-runtime/source/adapters/level_zero/v2/memory.cpp index 346432fbba54..3eaf09cf02ca 100644 --- a/unified-runtime/source/adapters/level_zero/v2/memory.cpp +++ b/unified-runtime/source/adapters/level_zero/v2/memory.cpp @@ -9,8 +9,10 @@ #include "memory.hpp" +#ifdef __linux__ #include #include +#endif #include "../ur_interface_loader.hpp" #include "context.hpp" @@ -933,6 +935,7 @@ ur_result_t urIPCGetPhysMemHandleExp(ur_context_handle_t hContext, ur_physical_mem_handle_t hPhysMem, void **ppIPCPhysMemHandleData, size_t *pIPCPhysMemHandleDataSizeRet) { +#ifdef __linux__ if (!hContext) return UR_RESULT_ERROR_INVALID_NULL_HANDLE; if (!hPhysMem) @@ -971,10 +974,18 @@ ur_result_t urIPCGetPhysMemHandleExp(ur_context_handle_t hContext, *ppIPCPhysMemHandleData = HandleData; *pIPCPhysMemHandleDataSizeRet = sizeof(ZeIPCPhysMemHandleData); return UR_RESULT_SUCCESS; +#else + (void)hContext; + (void)hPhysMem; + (void)ppIPCPhysMemHandleData; + (void)pIPCPhysMemHandleDataSizeRet; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +#endif // __linux__ } ur_result_t urIPCPutPhysMemHandleExp(ur_context_handle_t hContext, void *pIPCPhysMemHandleData) { +#ifdef __linux__ if (!hContext) return UR_RESULT_ERROR_INVALID_NULL_HANDLE; if (!pIPCPhysMemHandleData) @@ -985,6 +996,11 @@ ur_result_t urIPCPutPhysMemHandleExp(ur_context_handle_t hContext, close(HandleData->Fd); delete HandleData; return UR_RESULT_SUCCESS; +#else + (void)hContext; + (void)pIPCPhysMemHandleData; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +#endif // __linux__ } ur_result_t urIPCOpenPhysMemHandleExp(ur_context_handle_t hContext, @@ -992,6 +1008,7 @@ ur_result_t urIPCOpenPhysMemHandleExp(ur_context_handle_t hContext, void *pIPCPhysMemHandleData, size_t ipcPhysMemHandleDataSize, ur_physical_mem_handle_t *phPhysMem) { +#ifdef __linux__ if (!hContext) return UR_RESULT_ERROR_INVALID_NULL_HANDLE; if (!hDevice) @@ -1058,6 +1075,14 @@ ur_result_t urIPCOpenPhysMemHandleExp(ur_context_handle_t hContext, return UR_RESULT_ERROR_UNKNOWN; } return UR_RESULT_SUCCESS; +#else + (void)hContext; + (void)hDevice; + (void)pIPCPhysMemHandleData; + (void)ipcPhysMemHandleDataSize; + (void)phPhysMem; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +#endif // __linux__ } ur_result_t urIPCClosePhysMemHandleExp(ur_context_handle_t hContext, From 72e2b03d8748b2fea1ec3d97c513b4b9fde7f81c Mon Sep 17 00:00:00 2001 From: Lukasz Dorau Date: Mon, 25 May 2026 10:17:55 +0000 Subject: [PATCH 17/17] [SYCL] Fix MSVC C2487 and ABI dump for IPC physical_mem API Three fixes in this commit: 1. Remove __SYCL_EXPORT from private constructors in physical_mem.hpp. On MSVC, applying __declspec(dllexport) to a member of an already __declspec(dllexport) class causes C2487. The class-level __SYCL_EXPORT is sufficient to export all members on both Linux and Windows. 2. Add 3-param constructor (delegating to 4-param with false) to physical_mem.cpp. This constructor is declared in the header for ABI compatibility with existing compiled code. 3. Restore correct sycl_symbols_linux.dump header (was accidentally reordered) and add back 3-param constructor symbols (_ZN...Em) that were wrongly removed. Add new Windows symbols for: - physical_mem 4-param bool constructor - ipc::physical_memory::get/put/open - detail::openIPCPhysMemHandle - physical_mem::ext_oneapi_ipc_enabled Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- .../sycl/ext/oneapi/virtual_mem/physical_mem.hpp | 13 +++++++++---- sycl/source/physical_mem.cpp | 4 ++++ sycl/test/abi/sycl_symbols_linux.dump | 12 ++++++++---- sycl/test/abi/sycl_symbols_windows.dump | 6 ++++++ 4 files changed, 27 insertions(+), 8 deletions(-) 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 1cc1b39f5c4d..b87fa1c374b8 100644 --- a/sycl/include/sycl/ext/oneapi/virtual_mem/physical_mem.hpp +++ b/sycl/include/sycl/ext/oneapi/virtual_mem/physical_mem.hpp @@ -87,10 +87,15 @@ class __SYCL_EXPORT physical_mem bool ipc_enabled() const { return ext_oneapi_ipc_enabled(); } private: - // Internal constructor called by the public templated constructors. - __SYCL_EXPORT physical_mem(const device &SyclDevice, - const context &SyclContext, size_t NumBytes, - bool EnableIpc); + // Internal constructor called by the public templated constructors when IPC + // is not requested. Preserved for ABI compatibility. + physical_mem(const device &SyclDevice, const context &SyclContext, + size_t NumBytes); + + // Internal constructor called by the public templated constructors when IPC + // is requested. + physical_mem(const device &SyclDevice, const context &SyclContext, + size_t NumBytes, bool EnableIpc); // Internal constructor for creating a physical_mem from an existing impl // (used by createSyclObjFromImpl, e.g. when opening from an IPC handle). diff --git a/sycl/source/physical_mem.cpp b/sycl/source/physical_mem.cpp index c60ba3eed09c..868f16584b9c 100644 --- a/sycl/source/physical_mem.cpp +++ b/sycl/source/physical_mem.cpp @@ -13,6 +13,10 @@ namespace sycl { inline namespace _V1 { namespace ext::oneapi::experimental { +physical_mem::physical_mem(const device &SyclDevice, const context &SyclContext, + size_t NumBytes) + : physical_mem(SyclDevice, SyclContext, NumBytes, false) {} + physical_mem::physical_mem(const device &SyclDevice, const context &SyclContext, size_t NumBytes, bool EnableIpc) { if (!SyclDevice.has(aspect::ext_oneapi_virtual_mem)) diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index d93a8eea57c5..82f0d0b55e31 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -1,10 +1,12 @@ - +################################################################################ +# This file is automatically generated by abi_check.py tool. # DO NOT EDIT IT MANUALLY. Refer to sycl/doc/developer/ABIPolicyGuide.md for more info. -# REQUIRES: linux +################################################################################ + # RUN: env LLVM_BIN_PATH=%llvm_build_bin_dir %python %sycl_tools_src_dir/abi_check.py --mode check_symbols --reference %s %sycl_libs_dir/libsycl.so -# This file is automatically generated by abi_check.py tool. +# REQUIRES: linux # UNSUPPORTED: libcxx -################################################################################ + _ZN4sycl3_V110__abs_implENS0_3vecIaLi16EEE _ZN4sycl3_V110__abs_implENS0_3vecIaLi1EEE _ZN4sycl3_V110__abs_implENS0_3vecIaLi2EEE @@ -3004,7 +3006,9 @@ _ZN4sycl3_V13ext6oneapi12experimental12create_imageERNS3_9image_memERKNS3_16imag _ZN4sycl3_V13ext6oneapi12experimental12create_imageERNS3_9image_memERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental12create_imageERNS3_9image_memERKNS3_22bindless_image_samplerERKNS3_16image_descriptorERKNS0_5queueE _ZN4sycl3_V13ext6oneapi12experimental12create_imageERNS3_9image_memERKNS3_22bindless_image_samplerERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental12physical_memC1ERKNS0_6deviceERKNS0_7contextEm _ZN4sycl3_V13ext6oneapi12experimental12physical_memC1ERKNS0_6deviceERKNS0_7contextEmb +_ZN4sycl3_V13ext6oneapi12experimental12physical_memC2ERKNS0_6deviceERKNS0_7contextEm _ZN4sycl3_V13ext6oneapi12experimental12physical_memC2ERKNS0_6deviceERKNS0_7contextEmb _ZN4sycl3_V13ext6oneapi12experimental13aligned_allocEmmRKNS0_6deviceENS0_3usm5allocERKNS0_13property_listE _ZN4sycl3_V13ext6oneapi12experimental13malloc_deviceEmRKNS0_6deviceERKNS0_13property_listE diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index f07fa3e2e50b..d8b95d212f8d 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -418,6 +418,7 @@ ??0physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@$$QEAV012345@@Z ??0physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBV012345@@Z ??0physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBVdevice@45@AEBVcontext@45@_K@Z +??0physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBVdevice@45@AEBVcontext@45@_K_N@Z ??0physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBVqueue@45@_K@Z ??0platform@_V1@sycl@@AEAA@AEBVdevice@12@@Z ??0platform@_V1@sycl@@AEAA@V?$shared_ptr@Vplatform_impl@detail@_V1@sycl@@@std@@@Z @@ -3891,6 +3892,7 @@ ?ext_oneapi_has_kernel@kernel_bundle_plain@detail@_V1@sycl@@AEAA_NVstring_view@234@@Z ?ext_oneapi_has_kernel@kernel_bundle_plain@detail@_V1@sycl@@QEAA_NAEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@Z ?ext_oneapi_index_within_platform@device@_V1@sycl@@QEBA_KXZ +?ext_oneapi_ipc_enabled@physical_mem@experimental@oneapi@ext@_V1@sycl@@QEBA_NXZ ?ext_oneapi_memcpy2d_impl@handler@_V1@sycl@@AEAAXPEAX_KPEBX111@Z ?ext_oneapi_memset2d_impl@handler@_V1@sycl@@AEAAXPEAX_KH11@Z ?ext_oneapi_owner_before@?$OwnerLessBase@Vcontext@_V1@sycl@@@detail@_V1@sycl@@QEBA_NAEBV?$weak_object_base@Vcontext@_V1@sycl@@@2oneapi@ext@34@@Z @@ -3963,6 +3965,7 @@ ?get@ipc_memory@experimental@oneapi@ext@_V1@sycl@@YA?AUhandle@123456@PEAXAEBVcontext@56@@Z ?get@kernel@_V1@sycl@@QEBAPEAU_cl_kernel@@XZ ?get@memory@ipc@experimental@oneapi@ext@_V1@sycl@@YA?AUhandle@234567@PEAXAEBVcontext@67@@Z +?get@physical_memory@ipc@experimental@oneapi@ext@_V1@sycl@@YA?AUhandle@234567@AEBVphysical_mem@34567@@Z ?get@platform@_V1@sycl@@QEBAPEAU_cl_platform_id@@XZ ?get@queue@_V1@sycl@@QEBAPEAU_cl_command_queue@@XZ ?getAccData@AccessorBaseHost@detail@_V1@sycl@@QEAAAEAUAccHostDataT@234@XZ @@ -4261,7 +4264,9 @@ ?modf_impl@detail@_V1@sycl@@YANNPEAN@Z ?name@SYCLCategory@detail@_V1@sycl@@UEBAPEBDXZ ?native_specialization_constant@kernel_bundle_plain@detail@_V1@sycl@@QEBA_NXZ +?open@physical_memory@ipc@experimental@oneapi@ext@_V1@sycl@@YA?AVphysical_mem@34567@AEBV?$vector@W4byte@std@@V?$allocator@W4byte@std@@@2@@std@@AEBVcontext@67@AEBVdevice@67@@Z ?openIPCMemHandle@detail@_V1@sycl@@YAPEAXPEBW4byte@std@@_KAEBVcontext@23@AEBVdevice@23@@Z +?openIPCPhysMemHandle@detail@_V1@sycl@@YA?AVphysical_mem@experimental@oneapi@ext@23@PEBW4byte@std@@_KAEBVcontext@23@AEBVdevice@23@@Z ?parallel_for@handler@_V1@sycl@@QEAAXV?$range@$00@23@Vkernel@23@@Z ?parallel_for@handler@_V1@sycl@@QEAAXV?$range@$01@23@Vkernel@23@@Z ?parallel_for@handler@_V1@sycl@@QEAAXV?$range@$02@23@Vkernel@23@@Z @@ -4282,6 +4287,7 @@ ?print_graph@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@QEBAXV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@_N@Z ?put@ipc_memory@experimental@oneapi@ext@_V1@sycl@@YAXAEAUhandle@123456@AEBVcontext@56@@Z ?put@memory@ipc@experimental@oneapi@ext@_V1@sycl@@YAXAEAUhandle@234567@AEBVcontext@67@@Z +?put@physical_memory@ipc@experimental@oneapi@ext@_V1@sycl@@YAXAEAUhandle@234567@AEBVcontext@67@@Z ?query@tls_code_loc_t@detail@_V1@sycl@@QEAAAEBUcode_location@234@XZ ?reduComputeWGSize@detail@_V1@sycl@@YA_K_K0AEA_K@Z ?reduGetMaxNumConcurrentWorkGroups@detail@_V1@sycl@@YAIAEAVhandler@23@@Z