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/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 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..b87fa1c374b8 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,31 @@ 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 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). + 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/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/physical_mem_impl.hpp b/sycl/source/detail/physical_mem_impl.hpp index 3bc4d6865187..c71ab06f200a 100644 --- a/sycl/source/detail/physical_mem_impl.hpp +++ b/sycl/source/detail/physical_mem_impl.hpp @@ -38,13 +38,19 @@ 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(); + 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, Flags}; + 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 || @@ -55,9 +61,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 +96,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 +106,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/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/source/ipc_memory.cpp b/sycl/source/ipc_memory.cpp index 12b47b59d134..f59b32152777 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,58 @@ __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."); + + // 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 namespace ext::oneapi::experimental::ipc::memory { @@ -85,6 +138,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 +174,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..868f16584b9c 100644 --- a/sycl/source/physical_mem.cpp +++ b/sycl/source/physical_mem.cpp @@ -14,14 +14,23 @@ inline namespace _V1 { namespace ext::oneapi::experimental { physical_mem::physical_mem(const device &SyclDevice, const context &SyclContext, - size_t NumBytes) { + 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)) 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 +41,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/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/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 02ff46c47878..82f0d0b55e31 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3007,7 +3007,9 @@ _ZN4sycl3_V13ext6oneapi12experimental12create_imageERNS3_9image_memERKNS3_16imag _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 _ZN4sycl3_V13ext6oneapi12experimental13malloc_sharedEmRKNS0_6deviceERKNS0_13property_listE @@ -3076,6 +3078,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 +3324,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 +3654,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/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 diff --git a/unified-runtime/include/unified-runtime/ur_api.h b/unified-runtime/include/unified-runtime/ur_api.h index ce889941d39c..e18d7000cf2e 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 @@ -2503,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, @@ -5334,15 +5345,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. @@ -11240,6 +11251,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 +16577,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..16ab1442abb6 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; @@ -3306,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; @@ -5591,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) { @@ -9141,8 +9169,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"; @@ -9160,14 +9188,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); @@ -21694,6 +21723,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 +23385,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/scripts/core/exp-inter-process-communication.yml b/unified-runtime/scripts/core/exp-inter-process-communication.yml index e5019bb40836..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 @@ -126,3 +136,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 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 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/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/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/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/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 f0bdf7d23e59..bec36b889240 100644 --- a/unified-runtime/source/adapters/level_zero/memory.cpp +++ b/unified-runtime/source/adapters/level_zero/memory.cpp @@ -10,6 +10,10 @@ #include #include #include +#ifdef __linux__ +#include +#include +#endif #include #include "context.hpp" @@ -17,6 +21,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,6 +2016,173 @@ ur_result_t urIPCCloseMemHandleExp(ur_context_handle_t, void *pMem) { return umf::umf2urResult(umfCloseIPCHandle(pMem)); } +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) + 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 = {}; + 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)); + + auto *HandleData = new (std::nothrow) ZeIPCPhysMemHandleData; + if (!HandleData) { + close(ExportFd.fd); + 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; + + *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) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + + auto *HandleData = + static_cast(pIPCPhysMemHandleData); + 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, + ur_device_handle_t hDevice, + void *pIPCPhysMemHandleData, + size_t ipcPhysMemHandleDataSize, + ur_physical_mem_handle_t *phPhysMem) { +#ifdef __linux__ + 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); + + // 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 = ImportFdNum; + + 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); + // 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, hDevice, + HandleData->Size, + /*EnableIpc=*/false); + } 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; +#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, + ur_physical_mem_handle_t hPhysMem) { + if (!hContext) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + if (!hPhysMem) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + + // 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 // If indirect access tracking is enabled then performs reference counting, diff --git a/unified-runtime/source/adapters/level_zero/physical_mem.cpp b/unified-runtime/source/adapters/level_zero/physical_mem.cpp index a0663253f42e..563476415825 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, + hDevice, size, EnableIpc); } 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; @@ -66,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 a5db639575e8..290d6c826f46 100644 --- a/unified-runtime/source/adapters/level_zero/physical_mem.hpp +++ b/unified-runtime/source/adapters/level_zero/physical_mem.hpp @@ -11,10 +11,27 @@ #include "common.hpp" #include "common/ur_ref_count.hpp" +// Opaque handle data exchanged between processes for physical memory IPC. +// 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. +// 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, - ur_context_handle_t Context) - : ZePhysicalMem{ZePhysicalMem}, Context{Context} {} + 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; @@ -22,5 +39,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/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/level_zero/v2/memory.cpp b/unified-runtime/source/adapters/level_zero/v2/memory.cpp index b0601ba956af..3eaf09cf02ca 100644 --- a/unified-runtime/source/adapters/level_zero/v2/memory.cpp +++ b/unified-runtime/source/adapters/level_zero/v2/memory.cpp @@ -9,11 +9,17 @@ #include "memory.hpp" +#ifdef __linux__ +#include +#include +#endif + #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,4 +931,171 @@ ur_result_t urIPCCloseMemHandleExp(ur_context_handle_t, void *pMem) { return umf::umf2urResult(umfCloseIPCHandle(pMem)); } +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) + 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 = {}; + 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)); + + auto *HandleData = new (std::nothrow) ZeIPCPhysMemHandleData; + if (!HandleData) { + close(ExportFd.fd); + 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; + + *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) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + + auto *HandleData = + static_cast(pIPCPhysMemHandleData); + 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, + ur_device_handle_t hDevice, + void *pIPCPhysMemHandleData, + size_t ipcPhysMemHandleDataSize, + ur_physical_mem_handle_t *phPhysMem) { +#ifdef __linux__ + 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); + + // 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 = ImportFdNum; + + 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); + // 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, hDevice, + HandleData->Size, + /*EnableIpc=*/false); + } 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; +#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, + ur_physical_mem_handle_t hPhysMem) { + if (!hContext) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + if (!hPhysMem) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + + // 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/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/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/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/common.hpp b/unified-runtime/source/adapters/opencl/common.hpp index aefea1f6c55a..6644b42121cc 100644 --- a/unified-runtime/source/adapters/opencl/common.hpp +++ b/unified-runtime/source/adapters/opencl/common.hpp @@ -244,17 +244,17 @@ using clEnqueueReadGlobalVariableINTEL_fn = CL_API_ENTRY 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; 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( 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; } 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. 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..93430e0dacc8 --- /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_PHYSICAL_MEMORY_SUPPORT_EXP, + sizeof(ur_bool_t), &ipc_support, nullptr)); + if (!ipc_support) { + GTEST_SKIP() << "IPC physical 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; +} 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;