diff --git a/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td b/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td index b4f7c71af7d9f..9ad107e4874a6 100644 --- a/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td +++ b/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td @@ -97,6 +97,7 @@ def Aspectext_oneapi_exportable_device_mem : Aspect<"ext_oneapi_exportable_devic def Aspectext_oneapi_clock_sub_group : Aspect<"ext_oneapi_clock_sub_group">; 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_ipc_memory : Aspect<"ext_oneapi_ipc_memory">; // Deprecated aspects def AspectInt64_base_atomics : Aspect<"int64_base_atomics">; @@ -174,7 +175,8 @@ def : TargetInfo<"__TestAspectList", Aspectext_oneapi_exportable_device_mem, Aspectext_oneapi_clock_sub_group, Aspectext_oneapi_clock_work_group, - Aspectext_oneapi_clock_device], + Aspectext_oneapi_clock_device, + Aspectext_oneapi_ipc_memory], []>; // This definition serves the only purpose of testing whether the deprecated aspect list defined in here and in SYCL RT // match. 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 new file mode 100644 index 0000000000000..aa285570187d4 --- /dev/null +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_inter_process_communication.asciidoc @@ -0,0 +1,225 @@ += sycl_ext_oneapi_inter_process_communication + +:source-highlighter: coderay +:coderay-linenums-mode: table + +// This section needs to be after the document title. +:doctype: book +:toc2: +:toc: left +:encoding: utf-8 +:lang: en +:dpcpp: pass:[DPC++] +:endnote: —{nbsp}end{nbsp}note + +// Set the default source code type in this document to C++, +// for syntax highlighting purposes. This is needed because +// docbook uses c++ and html5 uses cpp. +:language: {basebackend@docbook:c++:cpp} + + +== Notice + +[%hardbreaks] +Copyright (C) 2025 Intel Corporation. All rights reserved. + +Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks +of The Khronos Group Inc. OpenCL(TM) is a trademark of Apple Inc. used by +permission by Khronos. + + +== Contact + +To report problems with this extension, please open a new issue at: + +https://github.com/intel/llvm/issues + + +== Dependencies + +This extension is written against the SYCL 2020 revision 10 specification. All +references below to the "core SYCL specification" or to section numbers in the +SYCL specification refer to that revision. + + +== Status + +This is an experimental extension specification, intended to provide early +access to features and gather community feedback. Interfaces defined in this +specification are implemented in {dpcpp}, but they are not finalized and may +change incompatibly in future versions of {dpcpp} without prior notice. +*Shipping software products should not rely on APIs defined in this +specification.* + + +== Backend support status + +The APIs in this extension may be used only on a device that has +`aspect::ext_oneapi_ipc_memory`. The application must check that the device has +this aspect before submitting a kernel using any of the APIs in this +extension. If the application fails to do this, the implementation throws +a synchronous exception with the `errc::kernel_not_supported` error code +when the kernel is submitted to the queue. + + +== 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` object and transferring the "handle data" to the other +processes. The other processes can use the handle data to recreate the +`ipc_memory` object and get a pointer to the corresponding device USM memory. + + +== Specification + +=== Feature test macro + +This extension provides a feature-test macro as described in the core SYCL +specification. An implementation supporting this extension must predefine the +macro `SYCL_EXT_ONEAPI_IPC` to one of the values defined in the table +below. Applications can test for the existence of this macro to determine if +the implementation supports this feature, or applications can test the macro's +value to determine which of the extension's features the implementation +supports. + +_And follow the text with a table like this *unless the extension is +"experimental"*. Note that your table may have more than one row if it +has multiple versions._ + +[%header,cols="1,5"] +|=== +|Value +|Description + +|1 +|The APIs of this experimental extension are not versioned, so the + feature-test macro always has this value. +|=== + +=== Inter-process communicable memory + + +This extension adds the new `ipc_memory` class. This new class adheres to the +common reference semantics described in +https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:reference-semantics[Section 4.5.2.] +in the SYCL 2020 specification. + +``` +namespace sycl::ext::oneapi::experimental { +using ipc_memory_handle_data_t = span; + +class ipc_memory { +public: + ipc_memory(void *ptr, sycl::context &ctx); + + static void *open(ipc_memory_handle_data_t ipc_memory_handle_data, + const sycl::context &ctx, const sycl::device &dev); + static void close(void *ptr, const sycl::context &ctx); + + ipc_memory_handle_data_t get_handle_data() const; + + void *get_ptr() const; +}; + +} +``` + +|==== +a| +[frame=all,grid=none] +!==== +a! +[source] +---- +ipc_memory(void *ptr, const sycl::context &ctx) +---- +!==== + +_Effects:_ Constructs an IPC memory object in `ctx` from a pointer `ptr` to +device USM memory. +If `ptr` is not pointing to device USM memory, the behaviors of this constructor +and any resulting objects are undefined. + +!==== +a! +[source] +---- +static void *open(ipc_memory_handle_data_t ipc_memory_handle_data, + const sycl::context &ctx, const sycl::device &dev) +---- +!==== + +_Effects:_ Returns a pointer to the same device USM memory as the device USM +memory associated with the `ipc_memory` object that the handle data originated +from. +The `ipc_memory` object that the handle data originated from is allowed to be +from another process on the host system. +If the `ipc_memory` object that the handle data originated from has been +destroyed, the behaviors of this constructor and any resulting objects are +undefined. +If the device USM memory the original `ipc_memory` object was created with was +not originally allocated on `dev`, the behaviors of this function is undefined. + +!==== +a! +[source] +---- +static void close(void *ptr, const sycl::context &ctx) +---- +!==== + +_Effects:_ Closes a device USM pointer previously returned by a call to +`ipc_memory::open()`. +Accessing `ptr` after a call to this function results in undefined behavior. + +!==== +a! +[source] +---- +ipc_memory_handle_data_t get_handle_data() const +---- +!==== + +_Returns:_ The handle data of the `ipc_memory` object. +Accessing the handle data returned by this API after the `ipc_memory` object has +been destroyed results in undefined behavior. + +!==== +a! +[source] +---- +void *get_ptr() const +---- +!==== + +_Returns:_ A pointer to device USM memory corresponding to the pointer used to +construct the original `ipc_memory` object. +Accessing the pointer returned by this API after the `ipc_memory` object has +been destroyed results in undefined behavior. + +|==== + + +== Issues + +=== Level Zero file descriptor duplication dependency + +The IPC memory APIs in Level Zero on Linux currently requires the ability to +duplicate file descriptors between processes. For security this is not allowed +by default on Linux-based systems, so in order for the IPC memory APIs to work +with Level Zero on Linux the user must either call `prctl(PR_SET_PTRACER, ...)` +in the IPC handle owner process or enable the functionality globally using + +```bash +sudo bash -c "echo 0 > /proc/sys/kernel/yama/ptrace_scope" +``` + +See also https://github.com/oneapi-src/unified-memory-framework/tree/main?tab=readme-ov-file#level-zero-memory-provider. + + +=== Level Zero IPC memory Windows support + +The new IPC memory APIs are not currently supported on the Level Zero backend on +Windows systems. + diff --git a/sycl/include/sycl/ext/oneapi/experimental/ipc_memory.hpp b/sycl/include/sycl/ext/oneapi/experimental/ipc_memory.hpp new file mode 100644 index 0000000000000..dbeb1aa9a6236 --- /dev/null +++ b/sycl/include/sycl/ext/oneapi/experimental/ipc_memory.hpp @@ -0,0 +1,63 @@ +//==------- ipc_memory.hpp --- SYCL inter-process communicable memory ------==// +// +// 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 +#include +#include +#include + +#include + +namespace sycl { +inline namespace _V1 { + +class context; +class device; + +namespace detail { +class ipc_memory_impl; +} + +namespace ext::oneapi::experimental { +using ipc_memory_handle_data_t = span; + +class __SYCL_EXPORT ipc_memory + : public sycl::detail::OwnerLessBase { +public: + ipc_memory(void *Ptr, const sycl::context &Ctx); + + static void *open(ipc_memory_handle_data_t IPCMemoryHandleData, + const sycl::context &Ctx, const sycl::device &Dev); + static void close(void *Ptr, const sycl::context &Ctx); + + ipc_memory_handle_data_t get_handle_data() const; + + void *get_ptr() const; + +private: + ipc_memory(std::shared_ptr IPCMemImpl) + : impl{IPCMemImpl} {} + + std::shared_ptr impl; + + template + friend const decltype(Obj::impl) & + sycl::detail::getSyclObjImpl(const Obj &SyclObject); + + template + friend T sycl::detail::createSyclObjFromImpl( + std::add_rvalue_reference_t ImplObj); + template + friend T sycl::detail::createSyclObjFromImpl( + std::add_lvalue_reference_t ImplObj); +}; +} // 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 d3e97a47a0248..f6c6b2e439bd4 100644 --- a/sycl/include/sycl/info/aspects.def +++ b/sycl/include/sycl/info/aspects.def @@ -83,3 +83,4 @@ __SYCL_ASPECT(ext_oneapi_exportable_device_mem, 90) __SYCL_ASPECT(ext_oneapi_clock_sub_group, 91) __SYCL_ASPECT(ext_oneapi_clock_work_group, 92) __SYCL_ASPECT(ext_oneapi_clock_device, 93) +__SYCL_ASPECT(ext_oneapi_ipc_memory, 94) diff --git a/sycl/include/sycl/sycl.hpp b/sycl/include/sycl/sycl.hpp index a09870dd77c30..e84dc848c42cc 100644 --- a/sycl/include/sycl/sycl.hpp +++ b/sycl/include/sycl/sycl.hpp @@ -128,6 +128,7 @@ can be disabled by setting SYCL_DISABLE_FSYCL_SYCLHPP_WARNING macro.") #include #include #include +#include #include #include #include diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt index bb11518b003b0..8272d3011dba3 100644 --- a/sycl/source/CMakeLists.txt +++ b/sycl/source/CMakeLists.txt @@ -321,6 +321,7 @@ set(SYCL_COMMON_SOURCES "handler.cpp" "image.cpp" "interop_handle.cpp" + "ipc_memory.cpp" "kernel.cpp" "kernel_bundle.cpp" "physical_mem.cpp" diff --git a/sycl/source/detail/device_impl.hpp b/sycl/source/detail/device_impl.hpp index 38214254595c6..f66a92d804ef4 100644 --- a/sycl/source/detail/device_impl.hpp +++ b/sycl/source/detail/device_impl.hpp @@ -1591,6 +1591,10 @@ class device_impl : public std::enable_shared_from_this { // Will be updated in a follow-up UR patch. return false; } + CASE(ext_oneapi_ipc_memory) { + return get_info_impl_nocheck() + .value_or(0); + } else { return false; // This device aspect has not been implemented yet. } diff --git a/sycl/source/detail/ipc_memory_impl.hpp b/sycl/source/detail/ipc_memory_impl.hpp new file mode 100644 index 0000000000000..b96cf2d3ce8b4 --- /dev/null +++ b/sycl/source/detail/ipc_memory_impl.hpp @@ -0,0 +1,78 @@ +//==-------- ipc_memory_impl.hpp --- SYCL ipc_memory implementation --------==// +// +// 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 +#include +#include +#include + +#include + +namespace sycl { +inline namespace _V1 { +namespace detail { + +class ipc_memory_impl { + struct private_tag { + explicit private_tag() = default; + }; + +public: + ipc_memory_impl(void *Ptr, const sycl::context &Ctx, private_tag) + : MContext{getSyclObjImpl(Ctx)}, MPtr{Ptr} { + adapter_impl &Adapter = MContext->getAdapter(); + Adapter.call(MContext->getHandleRef(), Ptr, + &MUrHandle); + } + + ipc_memory_impl(const ipc_memory_impl &) = delete; + ipc_memory_impl(ipc_memory_impl &&) = default; + + ~ipc_memory_impl() { + try { + adapter_impl &Adapter = MContext->getAdapter(); + Adapter.call_nocheck( + MContext->getHandleRef(), MUrHandle); + } catch (std::exception &e) { + __SYCL_REPORT_EXCEPTION_TO_STREAM("exception in ~ipc_memory_impl", e); + } + } + + ipc_memory_impl &operator=(const ipc_memory_impl &) = delete; + ipc_memory_impl &operator=(ipc_memory_impl &&) = default; + + template + static std::shared_ptr create(Ts &&...args) { + return std::make_shared(std::forward(args)..., + private_tag{}); + } + + sycl::ext::oneapi::experimental::ipc_memory_handle_data_t + get_handle_data() const { + adapter_impl &Adapter = MContext->getAdapter(); + void *HandleDataPtr = nullptr; + size_t HandleDataSize = 0; + Adapter.call( + MContext->getHandleRef(), MUrHandle, &HandleDataPtr, &HandleDataSize); + return sycl::span{ + reinterpret_cast(HandleDataPtr), HandleDataSize}; + } + + void *get_ptr() const { return MPtr; } + +private: + std::shared_ptr MContext; + void *MPtr = nullptr; + ur_exp_ipc_mem_handle_t MUrHandle = nullptr; +}; + +} // namespace detail +} // namespace _V1 +} // namespace sycl \ No newline at end of file diff --git a/sycl/source/detail/ur_device_info_ret_types.inc b/sycl/source/detail/ur_device_info_ret_types.inc index e1e724262b85f..31ef440822af7 100644 --- a/sycl/source/detail/ur_device_info_ret_types.inc +++ b/sycl/source/detail/ur_device_info_ret_types.inc @@ -162,6 +162,7 @@ MAP(UR_DEVICE_INFO_NODE_MASK, uint32_t) // These aren't present in the specification, extracted from ur_api.h // 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_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/feature_test.hpp.in b/sycl/source/feature_test.hpp.in index a641ae79b65e6..e8ca65def3a7c 100644 --- a/sycl/source/feature_test.hpp.in +++ b/sycl/source/feature_test.hpp.in @@ -126,6 +126,7 @@ inline namespace _V1 { #define SYCL_KHR_DEFAULT_CONTEXT 1 #define SYCL_EXT_INTEL_EVENT_MODE 1 #define SYCL_EXT_ONEAPI_TANGLE 1 +#define SYCL_EXT_ONEAPI_IPC 1 // Unfinished KHR extensions. These extensions are only available if the // __DPCPP_ENABLE_UNFINISHED_KHR_EXTENSIONS macro is defined. diff --git a/sycl/source/ipc_memory.cpp b/sycl/source/ipc_memory.cpp new file mode 100644 index 0000000000000..d6a8a81ba5fd5 --- /dev/null +++ b/sycl/source/ipc_memory.cpp @@ -0,0 +1,55 @@ +//==------- ipc_memory.cpp --- SYCL inter-process communicable memory ------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include +#include +#include +#include + +namespace sycl { +inline namespace _V1 { +namespace ext::oneapi::experimental { + +ipc_memory::ipc_memory(void *Ptr, const sycl::context &Ctx) + : impl(detail::ipc_memory_impl::create(Ptr, Ctx)) {} + +void *ipc_memory::open(ipc_memory_handle_data_t IPCMemoryHandleData, + const sycl::context &Ctx, const sycl::device &Dev) { + auto CtxImpl = sycl::detail::getSyclObjImpl(Ctx); + sycl::detail::adapter_impl &Adapter = CtxImpl->getAdapter(); + + void *Ptr = nullptr; + ur_result_t UrRes = + Adapter.call_nocheck( + CtxImpl->getHandleRef(), getSyclObjImpl(Dev)->getHandleRef(), + IPCMemoryHandleData.data(), IPCMemoryHandleData.size(), &Ptr); + if (UrRes == UR_RESULT_ERROR_INVALID_VALUE) + throw sycl::exception(sycl::make_error_code(errc::invalid), + "IPCMemoryHandleData data size does not correspond " + "to the target platform's IPC memory handle size."); + Adapter.checkUrResult(UrRes); + return Ptr; +} + +void ipc_memory::close(void *Ptr, const sycl::context &Ctx) { + auto CtxImpl = sycl::detail::getSyclObjImpl(Ctx); + sycl::detail::adapter_impl &Adapter = CtxImpl->getAdapter(); + Adapter.call( + CtxImpl->getHandleRef(), Ptr); +} + +ipc_memory_handle_data_t ipc_memory::get_handle_data() const { + return impl->get_handle_data(); +} + +void *ipc_memory::get_ptr() const { return impl->get_ptr(); } + +} // namespace ext::oneapi::experimental +} // namespace _V1 +} // namespace sycl diff --git a/sycl/test-e2e/Experimental/ipc_memory.cpp b/sycl/test-e2e/Experimental/ipc_memory.cpp new file mode 100644 index 0000000000000..ba14e874baa9b --- /dev/null +++ b/sycl/test-e2e/Experimental/ipc_memory.cpp @@ -0,0 +1,121 @@ +// REQUIRES: aspect-usm_device_allocations && aspect-ext_oneapi_ipc_memory + +// UNSUPPORTED: level_zero && windows +// UNSUPPORTED-TRACKER: UMFW-348 + +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +#include +#include +#include + +#include +#include +#include + +#if defined(__linux__) +#include +#include +#include +#endif // defined(__linux__) + +namespace syclexp = sycl::ext::oneapi::experimental; + +constexpr size_t N = 32; +constexpr const char *CommsFile = "ipc_comms.txt"; + +int spawner(int argc, char *argv[]) { + assert(argc == 1); + sycl::queue Q; + +#if defined(__linux__) + // UMF currently requires ptrace permissions to be set for the spawner. As + // such we need to set it until this limitation has been addressed. + // https://github.com/oneapi-src/unified-memory-framework/tree/main?tab=readme-ov-file#level-zero-memory-provider + if (Q.get_backend() == sycl::backend::ext_oneapi_level_zero && + prctl(PR_SET_PTRACER, getppid()) == -1) { + std::cout << "Failed to set ptracer permissions!" << std::endl; + return 1; + } +#endif // defined(__linux__) + + int *DataPtr = sycl::malloc_device(N, Q); + Q.parallel_for(N, [=](sycl::item<1> I) { + DataPtr[I] = static_cast(I.get_linear_id()); + }).wait(); + + { + syclexp::ipc_memory IPCMem{DataPtr, Q.get_context()}; + assert(IPCMem.get_ptr() == DataPtr); + + // Write handle data to file. + { + syclexp::ipc_memory_handle_data_t HandleData = IPCMem.get_handle_data(); + size_t HandleDataSize = HandleData.size(); + std::fstream FS(CommsFile, std::ios_base::out | std::ios_base::binary); + FS.write(reinterpret_cast(&HandleDataSize), sizeof(size_t)); + FS.write(HandleData.data(), HandleDataSize); + } + + // Spawn other process with an arguement. + std::string Cmd = std::string{argv[0]} + " 1"; + std::cout << "Spawning: " << Cmd << std::endl; + std::system(Cmd.c_str()); + } + + int Failures = 0; + int Read[N] = {0}; + Q.copy(DataPtr, Read, N).wait(); + for (size_t I = 0; I < N; ++I) { + if (Read[I] != (N - I)) { + ++Failures; + std::cout << "Failed from spawner: Result at " << I + << " unexpected: " << Read[I] << " != " << (N - I) << std::endl; + } + } + sycl::free(DataPtr, Q); + return Failures; +} + +int consumer() { + sycl::queue Q; + + // Read the handle data. + std::fstream FS(CommsFile, std::ios_base::in | std::ios_base::binary); + size_t HandleSize = 0; + FS.read(reinterpret_cast(&HandleSize), sizeof(size_t)); + std::unique_ptr HandleData{new char[HandleSize]}; + FS.read(HandleData.get(), HandleSize); + + // Open IPC handle. + syclexp::ipc_memory_handle_data_t Handle{HandleData.get(), HandleSize}; + int *DataPtr = reinterpret_cast( + syclexp::ipc_memory::open(Handle, Q.get_context(), Q.get_device())); + + // Test the data already in the USM pointer. + int Failures = 0; + int Read[N] = {0}; + Q.copy(DataPtr, Read, N).wait(); + for (size_t I = 0; I < N; ++I) { + if (Read[I] != I) { + ++Failures; + std::cout << "Failed from consumer: Result at " << I + << " unexpected: " << Read[I] << " != " << I << std::endl; + } + } + + Q.parallel_for(N, [=](sycl::item<1> I) { + DataPtr[I] = static_cast(N - I.get_linear_id()); + }).wait(); + + // Close the IPC pointer. + syclexp::ipc_memory::close(DataPtr, Q.get_context()); + + return Failures; +} + +int main(int argc, char *argv[]) { + // We either have the spawner (if no extra argument it provided) or + 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 e9f65ce662488..c64a07435c846 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -2991,6 +2991,10 @@ _ZN4sycl3_V13ext5intel12experimental9pipe_base18get_pipe_name_implEPKv _ZN4sycl3_V13ext6oneapi10level_zero6detail11make_deviceERKNS0_8platformEm _ZN4sycl3_V13ext6oneapi12experimental10async_freeERKNS0_5queueEPvRKNS0_6detail13code_locationE _ZN4sycl3_V13ext6oneapi12experimental10async_freeERNS0_7handlerEPv +_ZN4sycl3_V13ext6oneapi12experimental10ipc_memory4openENS0_4spanIcLm18446744073709551615EEERKNS0_7contextERKNS0_6deviceE +_ZN4sycl3_V13ext6oneapi12experimental10ipc_memory5closeEPvRKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental10ipc_memoryC1EPvRKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental10ipc_memoryC2EPvRKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental10mem_adviseENS0_5queueEPvmiRKNS0_6detail13code_locationE _ZN4sycl3_V13ext6oneapi12experimental11memory_pool21increase_threshold_toEm _ZN4sycl3_V13ext6oneapi12experimental11memory_poolC1ERKNS0_7contextERKNS0_6deviceENS0_3usm5allocENS4_15pool_propertiesE @@ -3691,6 +3695,8 @@ _ZNK4sycl3_V114interop_handle22ext_codeplay_has_graphEv _ZNK4sycl3_V115device_selector13select_deviceEv _ZNK4sycl3_V116default_selectorclERKNS0_6deviceE _ZNK4sycl3_V120accelerator_selectorclERKNS0_6deviceE +_ZNK4sycl3_V13ext6oneapi12experimental10ipc_memory7get_ptrEv +_ZNK4sycl3_V13ext6oneapi12experimental10ipc_memory15get_handle_dataEv _ZNK4sycl3_V13ext6oneapi12experimental11memory_pool10get_deviceEv _ZNK4sycl3_V13ext6oneapi12experimental11memory_pool11get_contextEv _ZNK4sycl3_V13ext6oneapi12experimental11memory_pool13get_thresholdEv diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index b3453854edec8..5bd0cf2657eb2 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -447,6 +447,10 @@ ??0image_plain@detail@_V1@sycl@@IEAA@_KAEBVcontext@23@Vevent@23@V?$unique_ptr@VSYCLMemObjAllocator@detail@_V1@sycl@@U?$default_delete@VSYCLMemObjAllocator@detail@_V1@sycl@@@std@@@std@@EW4image_channel_order@23@W4image_channel_type@23@_NV?$range@$02@23@@Z ??0image_plain@detail@_V1@sycl@@QEAA@$$QEAV0123@@Z ??0image_plain@detail@_V1@sycl@@QEAA@AEBV0123@@Z +??0ipc_memory@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBV012345@@Z +??0ipc_memory@experimental@oneapi@ext@_V1@sycl@@QEAA@PEAXAEBVcontext@45@@Z +??0ipc_memory@experimental@oneapi@ext@_V1@sycl@@AEAA@V?$shared_ptr@Vipc_memory_impl@detail@_V1@sycl@@@std@@@Z +??0ipc_memory@experimental@oneapi@ext@_V1@sycl@@QEAA@$$QEAV012345@@Z ??0kernel@_V1@sycl@@AEAA@V?$shared_ptr@Vkernel_impl@detail@_V1@sycl@@@std@@@Z ??0kernel@_V1@sycl@@QEAA@$$QEAV012@@Z ??0kernel@_V1@sycl@@QEAA@AEBV012@@Z @@ -537,6 +541,7 @@ ??1image_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@XZ ??1image_mem_impl@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@XZ ??1image_plain@detail@_V1@sycl@@QEAA@XZ +??1ipc_memory@experimental@oneapi@ext@_V1@sycl@@QEAA@XZ ??1kernel@_V1@sycl@@QEAA@XZ ??1kernel_bundle_plain@detail@_V1@sycl@@QEAA@XZ ??1kernel_id@_V1@sycl@@QEAA@XZ @@ -557,6 +562,8 @@ ??4?$OwnerLessBase@Vevent@_V1@sycl@@@detail@_V1@sycl@@QEAAAEAV0123@AEBV0123@@Z ??4?$OwnerLessBase@Vexecutable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@@detail@_V1@sycl@@QEAAAEAV0123@$$QEAV0123@@Z ??4?$OwnerLessBase@Vexecutable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@@detail@_V1@sycl@@QEAAAEAV0123@AEBV0123@@Z +??4?$OwnerLessBase@Vipc_memory@experimental@oneapi@ext@_V1@sycl@@@detail@_V1@sycl@@QEAAAEAV0123@AEBV0123@@Z +??4?$OwnerLessBase@Vipc_memory@experimental@oneapi@ext@_V1@sycl@@@detail@_V1@sycl@@QEAAAEAV0123@$$QEAV0123@@Z ??4?$OwnerLessBase@Vkernel@_V1@sycl@@@detail@_V1@sycl@@QEAAAEAV0123@$$QEAV0123@@Z ??4?$OwnerLessBase@Vkernel@_V1@sycl@@@detail@_V1@sycl@@QEAAAEAV0123@AEBV0123@@Z ??4?$OwnerLessBase@Vkernel_id@_V1@sycl@@@detail@_V1@sycl@@QEAAAEAV0123@$$QEAV0123@@Z @@ -635,6 +642,8 @@ ??4image_mem@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@AEBV012345@@Z ??4image_plain@detail@_V1@sycl@@QEAAAEAV0123@$$QEAV0123@@Z ??4image_plain@detail@_V1@sycl@@QEAAAEAV0123@AEBV0123@@Z +??4ipc_memory@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@AEBV012345@@Z +??4ipc_memory@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@$$QEAV012345@@Z ??4kernel@_V1@sycl@@QEAAAEAV012@$$QEAV012@@Z ??4kernel@_V1@sycl@@QEAAAEAV012@AEBV012@@Z ??4kernel_bundle_plain@detail@_V1@sycl@@QEAAAEAV0123@$$QEAV0123@@Z @@ -3830,6 +3839,7 @@ ?category@exception@_V1@sycl@@QEBAAEBVerror_category@std@@XZ ?checkNodePropertiesAndThrow@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@KAXAEBVproperty_list@67@@Z ?clearArgs@handler@_V1@sycl@@AEAAXXZ +?close@ipc_memory@experimental@oneapi@ext@_V1@sycl@@SAXPEAXAEBVcontext@56@@Z ?code@exception@_V1@sycl@@QEBAAEBVerror_code@std@@XZ ?compile_from_source@detail@experimental@oneapi@ext@_V1@sycl@@YA?AV?$kernel_bundle@$00@56@AEAV?$kernel_bundle@$02@56@AEBV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@std@@AEBV?$vector@Vstring_view@detail@_V1@sycl@@V?$allocator@Vstring_view@detail@_V1@sycl@@@std@@@std@@PEAVstring@156@2@Z ?compile_impl@detail@_V1@sycl@@YA?AV?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@AEBV?$kernel_bundle@$0A@@23@AEBV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@5@AEBVproperty_list@23@@Z @@ -3984,6 +3994,8 @@ ?ext_oneapi_owner_before@?$OwnerLessBase@Vevent@_V1@sycl@@@detail@_V1@sycl@@QEBA_NAEBVevent@34@@Z ?ext_oneapi_owner_before@?$OwnerLessBase@Vexecutable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@@detail@_V1@sycl@@QEBA_NAEBV?$weak_object_base@Vexecutable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@@2oneapi@ext@34@@Z ?ext_oneapi_owner_before@?$OwnerLessBase@Vexecutable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@@detail@_V1@sycl@@QEBA_NAEBVexecutable_command_graph@2experimental@oneapi@ext@34@@Z +?ext_oneapi_owner_before@?$OwnerLessBase@Vipc_memory@experimental@oneapi@ext@_V1@sycl@@@detail@_V1@sycl@@QEBA_NAEBV?$weak_object_base@Vipc_memory@experimental@oneapi@ext@_V1@sycl@@@2oneapi@ext@34@@Z +?ext_oneapi_owner_before@?$OwnerLessBase@Vipc_memory@experimental@oneapi@ext@_V1@sycl@@@detail@_V1@sycl@@QEBA_NAEBVipc_memory@experimental@oneapi@ext@34@@Z ?ext_oneapi_owner_before@?$OwnerLessBase@Vkernel@_V1@sycl@@@detail@_V1@sycl@@QEBA_NAEBV?$weak_object_base@Vkernel@_V1@sycl@@@2oneapi@ext@34@@Z ?ext_oneapi_owner_before@?$OwnerLessBase@Vkernel@_V1@sycl@@@detail@_V1@sycl@@QEBA_NAEBVkernel@34@@Z ?ext_oneapi_owner_before@?$OwnerLessBase@Vkernel_id@_V1@sycl@@@detail@_V1@sycl@@QEBA_NAEBV?$weak_object_base@Vkernel_id@_V1@sycl@@@2oneapi@ext@34@@Z @@ -4178,6 +4190,7 @@ ?get_filtering_mode@sampler@_V1@sycl@@QEBA?AW4filtering_mode@23@XZ ?get_flags@stream@_V1@sycl@@AEBAIXZ ?get_handle@image_mem@experimental@oneapi@ext@_V1@sycl@@QEBA?AUimage_mem_handle@23456@XZ +?get_handle_data@ipc_memory@experimental@oneapi@ext@_V1@sycl@@QEBA?AV?$span@D$0?0@56@XZ ?get_image_channel_type@experimental@oneapi@ext@_V1@sycl@@YA?AW4image_channel_type@45@Uimage_mem_handle@12345@AEBVdevice@45@AEBVcontext@45@@Z ?get_image_channel_type@experimental@oneapi@ext@_V1@sycl@@YA?AW4image_channel_type@45@Uimage_mem_handle@12345@AEBVqueue@45@@Z ?get_image_memory_support@experimental@oneapi@ext@_V1@sycl@@YA?AV?$vector@W4image_memory_handle_type@experimental@oneapi@ext@_V1@sycl@@V?$allocator@W4image_memory_handle_type@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@AEBUimage_descriptor@12345@AEBVdevice@45@AEBVcontext@45@@Z @@ -4217,6 +4230,7 @@ ?get_pointer_type@detail@_V1@sycl@@YA?AW4alloc@usm@23@PEBXAEAVcontext_impl@123@@Z ?get_precision@stream@_V1@sycl@@QEBA_KXZ ?get_predecessors@node@experimental@oneapi@ext@_V1@sycl@@QEBA?AV?$vector@Vnode@experimental@oneapi@ext@_V1@sycl@@V?$allocator@Vnode@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@XZ +?get_ptr@ipc_memory@experimental@oneapi@ext@_V1@sycl@@QEBAPEAXXZ ?get_queue@fusion_wrapper@experimental@codeplay@ext@_V1@sycl@@QEBA?AVqueue@56@XZ ?get_range@image_mem@experimental@oneapi@ext@_V1@sycl@@QEBA?AV?$range@$02@56@XZ ?get_range@image_plain@detail@_V1@sycl@@IEBA?AV?$range@$02@34@XZ @@ -4350,6 +4364,7 @@ ?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@ipc_memory@experimental@oneapi@ext@_V1@sycl@@SAPEAXV?$span@D$0?0@56@AEBVcontext@56@AEBVdevice@56@@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 diff --git a/sycl/unittests/Extensions/CMakeLists.txt b/sycl/unittests/Extensions/CMakeLists.txt index 59d57f0851ec1..3b3faf110bdbc 100644 --- a/sycl/unittests/Extensions/CMakeLists.txt +++ b/sycl/unittests/Extensions/CMakeLists.txt @@ -24,6 +24,7 @@ add_sycl_unittest(ExtensionsTests OBJECT DeviceInfo.cpp RootGroup.cpp USMPrefetch.cpp + IPC.cpp ) add_subdirectory(CommandGraph) diff --git a/sycl/unittests/Extensions/IPC.cpp b/sycl/unittests/Extensions/IPC.cpp new file mode 100644 index 0000000000000..6eff9c9429973 --- /dev/null +++ b/sycl/unittests/Extensions/IPC.cpp @@ -0,0 +1,163 @@ +//==------------------------------ IPC.cpp ---------------------------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include + +#include +#include +#include +#include +#include + +namespace syclexp = sycl::ext::oneapi::experimental; + +namespace { + +int DummyInt1 = 42; +void *DummyPtr = &DummyInt1; + +int DummyInt2 = 24; +ur_exp_ipc_mem_handle_t DummyMemHandle = + reinterpret_cast(&DummyInt2); + +constexpr size_t DummyHandleDataSize = 10; +char DummyHandleData[DummyHandleDataSize] = {9, 8, 7, 6, 5, 4, 3, 2, 1}; + +thread_local int urIPCGetMemHandleExp_counter = 0; +thread_local int urIPCPutMemHandleExp_counter = 0; +thread_local int urIPCOpenMemHandleExp_counter = 0; +thread_local int urIPCCloseMemHandleExp_counter = 0; +thread_local int urIPCGetMemHandleDataExp_counter = 0; + +ur_result_t replace_urIPCGetMemHandleExp(void *pParams) { + ++urIPCGetMemHandleExp_counter; + auto params = *static_cast(pParams); + EXPECT_EQ(*params.ppMem, DummyPtr); + **params.pphIPCMem = DummyMemHandle; + return UR_RESULT_SUCCESS; +} + +ur_result_t replace_urIPCPutMemHandleExp(void *pParams) { + ++urIPCPutMemHandleExp_counter; + auto params = *static_cast(pParams); + EXPECT_EQ(*params.phIPCMem, DummyMemHandle); + return UR_RESULT_SUCCESS; +} + +ur_result_t replace_urIPCOpenMemHandleExp(void *pParams) { + ++urIPCOpenMemHandleExp_counter; + auto params = *static_cast(pParams); + EXPECT_EQ(*params.pipcMemHandleData, DummyHandleData); + EXPECT_EQ(*params.pipcMemHandleDataSize, DummyHandleDataSize); + **params.pppMem = DummyPtr; + return UR_RESULT_SUCCESS; +} + +ur_result_t replace_urIPCCloseMemHandleExp(void *pParams) { + ++urIPCCloseMemHandleExp_counter; + auto params = *static_cast(pParams); + EXPECT_EQ(*params.ppMem, DummyPtr); + return UR_RESULT_SUCCESS; +} + +ur_result_t replace_urIPCGetMemHandleDataExp(void *pParams) { + ++urIPCGetMemHandleDataExp_counter; + auto params = + *static_cast(pParams); + EXPECT_EQ(*params.phIPCMem, DummyMemHandle); + **params.pppIPCHandleData = DummyHandleData; + **params.ppIPCMemHandleDataSizeRet = DummyHandleDataSize; + return UR_RESULT_SUCCESS; +} + +class IPCTests : public ::testing::Test { +public: + IPCTests() : Mock{}, Ctxt(sycl::platform()) {} + +protected: + void SetUp() override { + urIPCGetMemHandleExp_counter = 0; + urIPCPutMemHandleExp_counter = 0; + urIPCOpenMemHandleExp_counter = 0; + urIPCCloseMemHandleExp_counter = 0; + urIPCGetMemHandleDataExp_counter = 0; + + mock::getCallbacks().set_replace_callback("urIPCGetMemHandleExp", + replace_urIPCGetMemHandleExp); + mock::getCallbacks().set_replace_callback("urIPCPutMemHandleExp", + replace_urIPCPutMemHandleExp); + mock::getCallbacks().set_replace_callback("urIPCOpenMemHandleExp", + replace_urIPCOpenMemHandleExp); + mock::getCallbacks().set_replace_callback("urIPCCloseMemHandleExp", + replace_urIPCCloseMemHandleExp); + mock::getCallbacks().set_replace_callback("urIPCGetMemHandleDataExp", + replace_urIPCGetMemHandleDataExp); + } + + sycl::unittest::UrMock<> Mock; + sycl::context Ctxt; +}; + +TEST_F(IPCTests, IPCGetPut) { + { + syclexp::ipc_memory IPCMem{DummyPtr, Ctxt}; + + // Creating the IPC memory from a pointer should only call "get". + EXPECT_EQ(urIPCGetMemHandleExp_counter, 1); + EXPECT_EQ(urIPCPutMemHandleExp_counter, 0); + EXPECT_EQ(urIPCOpenMemHandleExp_counter, 0); + EXPECT_EQ(urIPCCloseMemHandleExp_counter, 0); + EXPECT_EQ(urIPCGetMemHandleDataExp_counter, 0); + + syclexp::ipc_memory_handle_data_t IPCMemHandleData = + IPCMem.get_handle_data(); + EXPECT_EQ(IPCMemHandleData.data(), DummyHandleData); + EXPECT_EQ(IPCMemHandleData.size(), DummyHandleDataSize); + + // Getting the underlying data should call the backend. + EXPECT_EQ(urIPCGetMemHandleExp_counter, 1); + EXPECT_EQ(urIPCPutMemHandleExp_counter, 0); + EXPECT_EQ(urIPCOpenMemHandleExp_counter, 0); + EXPECT_EQ(urIPCCloseMemHandleExp_counter, 0); + EXPECT_EQ(urIPCGetMemHandleDataExp_counter, 1); + } + + // When the IPC memory object dies, it should return the handle, calling + // "put". + EXPECT_EQ(urIPCGetMemHandleExp_counter, 1); + EXPECT_EQ(urIPCPutMemHandleExp_counter, 1); + EXPECT_EQ(urIPCOpenMemHandleExp_counter, 0); + EXPECT_EQ(urIPCCloseMemHandleExp_counter, 0); + EXPECT_EQ(urIPCGetMemHandleDataExp_counter, 1); +} + +TEST_F(IPCTests, IPCOpenClose) { + syclexp::ipc_memory_handle_data_t HandleData{DummyHandleData, + DummyHandleDataSize}; + void *Ptr = + syclexp::ipc_memory::open(HandleData, Ctxt, Ctxt.get_devices()[0]); + EXPECT_EQ(Ptr, DummyPtr); + + // Opening an IPC handle should call open. + EXPECT_EQ(urIPCGetMemHandleExp_counter, 0); + EXPECT_EQ(urIPCPutMemHandleExp_counter, 0); + EXPECT_EQ(urIPCOpenMemHandleExp_counter, 1); + EXPECT_EQ(urIPCCloseMemHandleExp_counter, 0); + EXPECT_EQ(urIPCGetMemHandleDataExp_counter, 0); + + syclexp::ipc_memory::close(Ptr, Ctxt); + + // When we close an IPC memory pointer, it should call close. + EXPECT_EQ(urIPCGetMemHandleExp_counter, 0); + EXPECT_EQ(urIPCPutMemHandleExp_counter, 0); + EXPECT_EQ(urIPCOpenMemHandleExp_counter, 1); + EXPECT_EQ(urIPCCloseMemHandleExp_counter, 1); + EXPECT_EQ(urIPCGetMemHandleDataExp_counter, 0); +} + +} // namespace diff --git a/unified-runtime/include/ur_api.h b/unified-runtime/include/ur_api.h index f78714b4e06aa..5602bc0b5d0e0 100644 --- a/unified-runtime/include/ur_api.h +++ b/unified-runtime/include/ur_api.h @@ -475,6 +475,16 @@ typedef enum ur_function_t { UR_FUNCTION_MEMORY_EXPORT_EXPORT_MEMORY_HANDLE_EXP = 287, /// Enumerator for ::urBindlessImagesSupportsImportingHandleTypeExp UR_FUNCTION_BINDLESS_IMAGES_SUPPORTS_IMPORTING_HANDLE_TYPE_EXP = 288, + /// Enumerator for ::urIPCGetMemHandleExp + UR_FUNCTION_IPC_GET_MEM_HANDLE_EXP = 289, + /// Enumerator for ::urIPCPutMemHandleExp + UR_FUNCTION_IPC_PUT_MEM_HANDLE_EXP = 290, + /// Enumerator for ::urIPCOpenMemHandleExp + UR_FUNCTION_IPC_OPEN_MEM_HANDLE_EXP = 291, + /// Enumerator for ::urIPCCloseMemHandleExp + UR_FUNCTION_IPC_CLOSE_MEM_HANDLE_EXP = 292, + /// Enumerator for ::urIPCGetMemHandleDataExp + UR_FUNCTION_IPC_GET_MEM_HANDLE_DATA_EXP = 293, /// @cond UR_FUNCTION_FORCE_UINT32 = 0x7fffffff /// @endcond @@ -2426,6 +2436,9 @@ typedef enum ur_device_info_t { /// [::ur_exp_device_2d_block_array_capability_flags_t] return a bit-field /// of Intel GPU 2D block array capabilities UR_DEVICE_INFO_2D_BLOCK_ARRAY_CAPABILITIES_EXP = 0x2022, + /// [::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 enqueueing of /// allocations and frees. UR_DEVICE_INFO_ASYNC_USM_ALLOCATIONS_SUPPORT_EXP = 0x2050, @@ -12347,6 +12360,141 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueTimestampRecordingExp( /// array. ur_event_handle_t *phEvent); +#if !defined(__GNUC__) +#pragma endregion +#endif +// Intel 'oneAPI' Unified Runtime Experimental APIs for Inter Process +// Communication +#if !defined(__GNUC__) +#pragma region inter_process_communication_(experimental) +#endif +/////////////////////////////////////////////////////////////////////////////// +/// @brief Handle of inter-process communicable memory +typedef struct ur_exp_ipc_mem_handle_t_ *ur_exp_ipc_mem_handle_t; + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Gets an inter-process memory handle for a pointer to device USM +/// memory +/// +/// @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 == phIPCMem` +/// - ::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 urIPCGetMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] pointer to device USM memory + void *pMem, + /// [out][alloc] pointer to the resulting IPC memory handle + ur_exp_ipc_mem_handle_t *phIPCMem); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Releases an inter-process 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 == hIPCMem` +/// - ::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 urIPCPutMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] the IPC memory handle + ur_exp_ipc_mem_handle_t hIPCMem); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Opens an inter-process memory handle from raw data to get the +/// corresponding pointer to device USM memory +/// +/// @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_CONTEXT +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == ipcMemHandleData` +/// + `NULL == ppMem` +/// - ::UR_RESULT_ERROR_INVALID_VALUE +/// + ipcMemHandleDataSize is not the same as the size of IPC memory +/// handle data +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +UR_APIEXPORT ur_result_t UR_APICALL urIPCOpenMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object the corresponding USM device memory + /// was allocated on + ur_device_handle_t hDevice, + /// [in] the IPC memory handle data + void *ipcMemHandleData, + /// [in] size of the IPC memory handle data + size_t ipcMemHandleDataSize, + /// [out] pointer to a pointer to device USM memory + void **ppMem); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Closes an inter-process 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_CONTEXT +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pMem` +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +UR_APIEXPORT ur_result_t UR_APICALL urIPCCloseMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] pointer to device USM memory opened through urIPCOpenMemHandleExp + void *pMem); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Gets the data of an inter-process 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 == hIPCMem` +/// - ::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 urIPCGetMemHandleDataExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] the IPC memory handle + ur_exp_ipc_mem_handle_t hIPCMem, + /// [out][optional] a pointer to the IPC memory handle data + void **ppIPCHandleData, + /// [out][optional] size of the resulting IPC memory handle data + size_t *pIPCMemHandleDataSizeRet); + #if !defined(__GNUC__) #pragma endregion #endif @@ -15422,6 +15570,57 @@ typedef struct ur_command_buffer_get_native_handle_exp_params_t { ur_native_handle_t **pphNativeCommandBuffer; } ur_command_buffer_get_native_handle_exp_params_t; +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function parameters for urIPCGetMemHandleExp +/// @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_mem_handle_exp_params_t { + ur_context_handle_t *phContext; + void **ppMem; + ur_exp_ipc_mem_handle_t **pphIPCMem; +} ur_ipc_get_mem_handle_exp_params_t; + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function parameters for urIPCPutMemHandleExp +/// @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_mem_handle_exp_params_t { + ur_context_handle_t *phContext; + ur_exp_ipc_mem_handle_t *phIPCMem; +} ur_ipc_put_mem_handle_exp_params_t; + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function parameters for urIPCOpenMemHandleExp +/// @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_mem_handle_exp_params_t { + ur_context_handle_t *phContext; + ur_device_handle_t *phDevice; + void **pipcMemHandleData; + size_t *pipcMemHandleDataSize; + void ***pppMem; +} ur_ipc_open_mem_handle_exp_params_t; + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function parameters for urIPCCloseMemHandleExp +/// @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_mem_handle_exp_params_t { + ur_context_handle_t *phContext; + void **ppMem; +} ur_ipc_close_mem_handle_exp_params_t; + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function parameters for urIPCGetMemHandleDataExp +/// @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_mem_handle_data_exp_params_t { + ur_context_handle_t *phContext; + ur_exp_ipc_mem_handle_t *phIPCMem; + void ***pppIPCHandleData; + size_t **ppIPCMemHandleDataSizeRet; +} ur_ipc_get_mem_handle_data_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/ur_api_funcs.def b/unified-runtime/include/ur_api_funcs.def index f0c92445b9238..4b72a768d8ca3 100644 --- a/unified-runtime/include/ur_api_funcs.def +++ b/unified-runtime/include/ur_api_funcs.def @@ -206,6 +206,11 @@ _UR_API(urCommandBufferUpdateSignalEventExp) _UR_API(urCommandBufferUpdateWaitEventsExp) _UR_API(urCommandBufferGetInfoExp) _UR_API(urCommandBufferGetNativeHandleExp) +_UR_API(urIPCGetMemHandleExp) +_UR_API(urIPCPutMemHandleExp) +_UR_API(urIPCOpenMemHandleExp) +_UR_API(urIPCCloseMemHandleExp) +_UR_API(urIPCGetMemHandleDataExp) _UR_API(urMemoryExportAllocExportableMemoryExp) _UR_API(urMemoryExportFreeExportableMemoryExp) _UR_API(urMemoryExportExportMemoryHandleExp) diff --git a/unified-runtime/include/ur_ddi.h b/unified-runtime/include/ur_ddi.h index 8ab686aa583cc..035be44791ab7 100644 --- a/unified-runtime/include/ur_ddi.h +++ b/unified-runtime/include/ur_ddi.h @@ -1787,6 +1787,61 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetCommandBufferExpProcAddrTable( typedef ur_result_t(UR_APICALL *ur_pfnGetCommandBufferExpProcAddrTable_t)( ur_api_version_t, ur_command_buffer_exp_dditable_t *); +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function-pointer for urIPCGetMemHandleExp +typedef ur_result_t(UR_APICALL *ur_pfnIPCGetMemHandleExp_t)( + ur_context_handle_t, void *, ur_exp_ipc_mem_handle_t *); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function-pointer for urIPCPutMemHandleExp +typedef ur_result_t(UR_APICALL *ur_pfnIPCPutMemHandleExp_t)( + ur_context_handle_t, ur_exp_ipc_mem_handle_t); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function-pointer for urIPCOpenMemHandleExp +typedef ur_result_t(UR_APICALL *ur_pfnIPCOpenMemHandleExp_t)( + ur_context_handle_t, ur_device_handle_t, void *, size_t, void **); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function-pointer for urIPCCloseMemHandleExp +typedef ur_result_t(UR_APICALL *ur_pfnIPCCloseMemHandleExp_t)( + ur_context_handle_t, void *); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function-pointer for urIPCGetMemHandleDataExp +typedef ur_result_t(UR_APICALL *ur_pfnIPCGetMemHandleDataExp_t)( + ur_context_handle_t, ur_exp_ipc_mem_handle_t, void **, size_t *); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Table of IPCExp functions pointers +typedef struct ur_ipc_exp_dditable_t { + ur_pfnIPCGetMemHandleExp_t pfnGetMemHandleExp; + ur_pfnIPCPutMemHandleExp_t pfnPutMemHandleExp; + ur_pfnIPCOpenMemHandleExp_t pfnOpenMemHandleExp; + ur_pfnIPCCloseMemHandleExp_t pfnCloseMemHandleExp; + ur_pfnIPCGetMemHandleDataExp_t pfnGetMemHandleDataExp; +} ur_ipc_exp_dditable_t; + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Exported function for filling application's IPCExp table +/// with current process' addresses +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// - ::UR_RESULT_ERROR_UNSUPPORTED_VERSION +UR_DLLEXPORT ur_result_t UR_APICALL urGetIPCExpProcAddrTable( + /// [in] API version requested + ur_api_version_t version, + /// [in,out] pointer to table of DDI function pointers + ur_ipc_exp_dditable_t *pDdiTable); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function-pointer for urGetIPCExpProcAddrTable +typedef ur_result_t(UR_APICALL *ur_pfnGetIPCExpProcAddrTable_t)( + ur_api_version_t, ur_ipc_exp_dditable_t *); + /////////////////////////////////////////////////////////////////////////////// /// @brief Function-pointer for urMemoryExportAllocExportableMemoryExp typedef ur_result_t(UR_APICALL *ur_pfnMemoryExportAllocExportableMemoryExp_t)( @@ -2049,6 +2104,7 @@ typedef struct ur_dditable_t { ur_usm_exp_dditable_t USMExp; ur_bindless_images_exp_dditable_t BindlessImagesExp; ur_command_buffer_exp_dditable_t CommandBufferExp; + ur_ipc_exp_dditable_t IPCExp; ur_memory_export_exp_dditable_t MemoryExportExp; ur_usm_p2p_exp_dditable_t UsmP2PExp; ur_virtual_mem_dditable_t VirtualMem; diff --git a/unified-runtime/include/ur_print.h b/unified-runtime/include/ur_print.h index 8130df0c5bec4..dd38252f71399 100644 --- a/unified-runtime/include/ur_print.h +++ b/unified-runtime/include/ur_print.h @@ -3490,6 +3490,56 @@ urPrintCommandBufferGetNativeHandleExpParams( const struct ur_command_buffer_get_native_handle_exp_params_t *params, char *buffer, const size_t buff_size, size_t *out_size); +/////////////////////////////////////////////////////////////////////////////// +/// @brief Print ur_ipc_get_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 urPrintIpcGetMemHandleExpParams( + const struct ur_ipc_get_mem_handle_exp_params_t *params, char *buffer, + const size_t buff_size, size_t *out_size); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Print ur_ipc_put_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 urPrintIpcPutMemHandleExpParams( + const struct ur_ipc_put_mem_handle_exp_params_t *params, char *buffer, + const size_t buff_size, size_t *out_size); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Print ur_ipc_open_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 urPrintIpcOpenMemHandleExpParams( + const struct ur_ipc_open_mem_handle_exp_params_t *params, char *buffer, + const size_t buff_size, size_t *out_size); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Print ur_ipc_close_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 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_mem_handle_data_exp_params_t struct +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_INVALID_SIZE +/// - `buff_size < out_size` +UR_APIEXPORT ur_result_t UR_APICALL urPrintIpcGetMemHandleDataExpParams( + const struct ur_ipc_get_mem_handle_data_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/ur_print.hpp b/unified-runtime/include/ur_print.hpp index 15c50dd0eb479..73448963a1dc3 100644 --- a/unified-runtime/include/ur_print.hpp +++ b/unified-runtime/include/ur_print.hpp @@ -42,6 +42,7 @@ template <> struct is_handle : std::true_type {}; template <> struct is_handle : std::true_type {}; +template <> struct is_handle : std::true_type {}; template inline constexpr bool is_handle_v = is_handle::value; template inline ur_result_t printPtr(std::ostream &os, const T *ptr); @@ -1274,6 +1275,21 @@ inline std::ostream &operator<<(std::ostream &os, enum ur_function_t value) { case UR_FUNCTION_BINDLESS_IMAGES_SUPPORTS_IMPORTING_HANDLE_TYPE_EXP: os << "UR_FUNCTION_BINDLESS_IMAGES_SUPPORTS_IMPORTING_HANDLE_TYPE_EXP"; break; + case UR_FUNCTION_IPC_GET_MEM_HANDLE_EXP: + os << "UR_FUNCTION_IPC_GET_MEM_HANDLE_EXP"; + break; + case UR_FUNCTION_IPC_PUT_MEM_HANDLE_EXP: + os << "UR_FUNCTION_IPC_PUT_MEM_HANDLE_EXP"; + break; + case UR_FUNCTION_IPC_OPEN_MEM_HANDLE_EXP: + os << "UR_FUNCTION_IPC_OPEN_MEM_HANDLE_EXP"; + break; + case UR_FUNCTION_IPC_CLOSE_MEM_HANDLE_EXP: + os << "UR_FUNCTION_IPC_CLOSE_MEM_HANDLE_EXP"; + break; + case UR_FUNCTION_IPC_GET_MEM_HANDLE_DATA_EXP: + os << "UR_FUNCTION_IPC_GET_MEM_HANDLE_DATA_EXP"; + break; default: os << "unknown enumerator"; break; @@ -3116,6 +3132,9 @@ inline std::ostream &operator<<(std::ostream &os, enum ur_device_info_t value) { case UR_DEVICE_INFO_2D_BLOCK_ARRAY_CAPABILITIES_EXP: os << "UR_DEVICE_INFO_2D_BLOCK_ARRAY_CAPABILITIES_EXP"; break; + case UR_DEVICE_INFO_IPC_MEMORY_SUPPORT_EXP: + os << "UR_DEVICE_INFO_IPC_MEMORY_SUPPORT_EXP"; + break; case UR_DEVICE_INFO_ASYNC_USM_ALLOCATIONS_SUPPORT_EXP: os << "UR_DEVICE_INFO_ASYNC_USM_ALLOCATIONS_SUPPORT_EXP"; break; @@ -5244,6 +5263,19 @@ inline ur_result_t printTagged(std::ostream &os, const void *ptr, os << ")"; } break; + case UR_DEVICE_INFO_IPC_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) { @@ -20258,6 +20290,138 @@ operator<<(std::ostream &os, [[maybe_unused]] const struct return os; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Print operator for the ur_ipc_get_mem_handle_exp_params_t type +/// @returns +/// std::ostream & +inline std::ostream &operator<<( + std::ostream &os, + [[maybe_unused]] const struct ur_ipc_get_mem_handle_exp_params_t *params) { + + os << ".hContext = "; + + ur::details::printPtr(os, *(params->phContext)); + + os << ", "; + os << ".pMem = "; + + os << *(params->ppMem); + + os << ", "; + os << ".phIPCMem = "; + + ur::details::printPtr(os, *(params->pphIPCMem)); + + return os; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Print operator for the ur_ipc_put_mem_handle_exp_params_t type +/// @returns +/// std::ostream & +inline std::ostream &operator<<( + std::ostream &os, + [[maybe_unused]] const struct ur_ipc_put_mem_handle_exp_params_t *params) { + + os << ".hContext = "; + + ur::details::printPtr(os, *(params->phContext)); + + os << ", "; + os << ".hIPCMem = "; + + ur::details::printPtr(os, *(params->phIPCMem)); + + return os; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Print operator for the ur_ipc_open_mem_handle_exp_params_t type +/// @returns +/// std::ostream & +inline std::ostream &operator<<( + std::ostream &os, + [[maybe_unused]] const struct ur_ipc_open_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 << ".ipcMemHandleData = "; + + os << *(params->pipcMemHandleData); + + os << ", "; + os << ".ipcMemHandleDataSize = "; + + os << *(params->pipcMemHandleDataSize); + + os << ", "; + os << ".ppMem = "; + + os << *(params->pppMem); + + return os; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Print operator for the ur_ipc_close_mem_handle_exp_params_t type +/// @returns +/// std::ostream & +inline std::ostream & +operator<<(std::ostream &os, + [[maybe_unused]] const struct ur_ipc_close_mem_handle_exp_params_t + *params) { + + os << ".hContext = "; + + ur::details::printPtr(os, *(params->phContext)); + + os << ", "; + os << ".pMem = "; + + os << *(params->ppMem); + + return os; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Print operator for the ur_ipc_get_mem_handle_data_exp_params_t type +/// @returns +/// std::ostream & +inline std::ostream & +operator<<(std::ostream &os, + [[maybe_unused]] const struct ur_ipc_get_mem_handle_data_exp_params_t + *params) { + + os << ".hContext = "; + + ur::details::printPtr(os, *(params->phContext)); + + os << ", "; + os << ".hIPCMem = "; + + ur::details::printPtr(os, *(params->phIPCMem)); + + os << ", "; + os << ".ppIPCHandleData = "; + + ur::details::printPtr(os, *(params->pppIPCHandleData)); + + os << ", "; + os << ".pIPCMemHandleDataSizeRet = "; + + ur::details::printPtr(os, *(params->ppIPCMemHandleDataSizeRet)); + + return os; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Print operator for the /// ur_memory_export_alloc_exportable_memory_exp_params_t type @@ -21722,6 +21886,21 @@ inline ur_result_t UR_APICALL printFunctionParams(std::ostream &os, os << (const struct ur_command_buffer_get_native_handle_exp_params_t *) params; } break; + case UR_FUNCTION_IPC_GET_MEM_HANDLE_EXP: { + os << (const struct ur_ipc_get_mem_handle_exp_params_t *)params; + } break; + case UR_FUNCTION_IPC_PUT_MEM_HANDLE_EXP: { + os << (const struct ur_ipc_put_mem_handle_exp_params_t *)params; + } break; + case UR_FUNCTION_IPC_OPEN_MEM_HANDLE_EXP: { + os << (const struct ur_ipc_open_mem_handle_exp_params_t *)params; + } break; + 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_MEM_HANDLE_DATA_EXP: { + os << (const struct ur_ipc_get_mem_handle_data_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.rst b/unified-runtime/scripts/core/EXP-INTER-PROCESS-COMMUNICATION.rst new file mode 100644 index 0000000000000..73d0616a1d7c3 --- /dev/null +++ b/unified-runtime/scripts/core/EXP-INTER-PROCESS-COMMUNICATION.rst @@ -0,0 +1,69 @@ +<% + OneApi=tags['$OneApi'] + x=tags['$x'] + X=x.upper() +%> + +.. _experimental-inter-process-communication: + +================================================================================ +Inter Process Communication +================================================================================ + +.. warning:: + + Experimental features: + + * May be replaced, updated, or removed at any time. + * Do not require maintaining API/ABI stability of their own additions over + time. + * Do not require conformance testing of their own additions. + + +Motivation +-------------------------------------------------------------------------------- +This extension introduces functionality for allowing processes to share common +objects, such as device USM memory allocations. Doing so lets processes actively +communicate with each other through the devices, by explicitly managing handles +that represent shareable objects for inter-process communication. + +API +-------------------------------------------------------------------------------- + +Enums +~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ +* ${x}_device_info_t + * ${X}_DEVICE_INFO_IPC_MEMORY_SUPPORT_EXP + +Types +~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ +* ${x}_exp_ipc_mem_handle_t + +Functions +~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ +* Inter-Process Communication + * ${x}IPCGetMemHandleExp + * ${x}IPCPutMemHandleExp + * ${x}IPCOpenMemHandleExp + * ${x}IPCCloseMemHandleExp + * ${x}IPCGetMemHandleDataExp + +Changelog +-------------------------------------------------------------------------------- + ++-----------+------------------------+ +| Revision | Changes | ++===========+========================+ +| 1.0 | Initial Draft | ++-----------+------------------------+ + +Support +-------------------------------------------------------------------------------- + +Adapters which support this experimental feature *must* return true for the new +``${X}_DEVICE_INFO_IPC_MEMORY_SUPPORT_EXP`` device info query. + +Contributors +-------------------------------------------------------------------------------- + +* Larsen, Steffen `steffen.larsen@intel.com `_ diff --git a/unified-runtime/scripts/core/exp-inter-process-communication.yml b/unified-runtime/scripts/core/exp-inter-process-communication.yml new file mode 100644 index 0000000000000..e5fb1691fe354 --- /dev/null +++ b/unified-runtime/scripts/core/exp-inter-process-communication.yml @@ -0,0 +1,154 @@ +# +# Copyright (C) 2025 Intel Corporation +# +# Part of the Unified-Runtime Project, under the Apache License v2.0 with LLVM +# Exceptions. +# See LICENSE.TXT +# +# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +# +# See YaML.md for syntax definition +# +--- #-------------------------------------------------------------------------- +type: header +desc: "Intel $OneApi Unified Runtime Experimental APIs for Inter Process Communication" +ordinal: "99" +--- #-------------------------------------------------------------------------- +type: enum +extend: true +typed_etors: true +desc: "Extension enums to $x_device_info_t to support inter-process communicable memory handles." +name: $x_device_info_t +etors: + - name: IPC_MEMORY_SUPPORT_EXP + value: "0x2023" + desc: "[$x_bool_t] returns true if the device supports inter-process communicable memory handles" +--- #-------------------------------------------------------------------------- +type: handle +desc: "Handle of inter-process communicable memory" +name: "$x_exp_ipc_mem_handle_t" +--- #-------------------------------------------------------------------------- +type: function +desc: "Gets an inter-process memory handle for a pointer to device USM memory" +class: $xIPC +name: GetMemHandleExp +ordinal: "0" +params: + - type: $x_context_handle_t + name: hContext + desc: "[in] handle of the context object" + - type: void * + name: pMem + desc: "[in] pointer to device USM memory" + - type: $x_exp_ipc_mem_handle_t* + name: phIPCMem + desc: "[out][alloc] pointer to the resulting IPC memory handle" +returns: + - $X_RESULT_ERROR_INVALID_CONTEXT + - $X_RESULT_ERROR_INVALID_NULL_HANDLE: + - "`NULL == hContext`" + - $X_RESULT_ERROR_INVALID_NULL_POINTER: + - "`NULL == phIPCMem`" + - $X_RESULT_ERROR_OUT_OF_HOST_MEMORY + - $X_RESULT_ERROR_OUT_OF_RESOURCES +--- #-------------------------------------------------------------------------- +type: function +desc: "Releases an inter-process memory handle" +class: $xIPC +name: PutMemHandleExp +ordinal: "0" +params: + - type: $x_context_handle_t + name: hContext + desc: "[in] handle of the context object" + - type: $x_exp_ipc_mem_handle_t + name: hIPCMem + desc: "[in] the IPC memory handle" +returns: + - $X_RESULT_ERROR_INVALID_CONTEXT + - $X_RESULT_ERROR_INVALID_NULL_HANDLE: + - "`NULL == hContext`" + - "`NULL == hIPCMem`" + - $X_RESULT_ERROR_OUT_OF_HOST_MEMORY + - $X_RESULT_ERROR_OUT_OF_RESOURCES +--- #-------------------------------------------------------------------------- +type: function +desc: "Opens an inter-process memory handle from raw data to get the corresponding pointer to device USM memory" +class: $xIPC +name: OpenMemHandleExp +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 corresponding USM device memory was allocated on" + - type: void * + name: ipcMemHandleData + desc: "[in] the IPC memory handle data" + - type: size_t + name: ipcMemHandleDataSize + desc: "[in] size of the IPC memory handle data" + - type: void ** + name: ppMem + desc: "[out] pointer to a pointer to device USM memory" +returns: + - $X_RESULT_ERROR_INVALID_CONTEXT + - $X_RESULT_ERROR_INVALID_NULL_HANDLE: + - "`NULL == hContext`" + - "`NULL == hDevice`" + - $X_RESULT_ERROR_INVALID_NULL_POINTER: + - "`NULL == ipcMemHandleData`" + - "`NULL == ppMem`" + - $X_RESULT_ERROR_INVALID_VALUE: + - "ipcMemHandleDataSize is not the same as the size of IPC memory handle data" + - $X_RESULT_ERROR_OUT_OF_HOST_MEMORY + - $X_RESULT_ERROR_OUT_OF_RESOURCES +--- #-------------------------------------------------------------------------- +type: function +desc: "Closes an inter-process memory handle" +class: $xIPC +name: CloseMemHandleExp +ordinal: "0" +params: + - type: $x_context_handle_t + name: hContext + desc: "[in] handle of the context object" + - type: void * + name: pMem + desc: "[in] pointer to device USM memory opened through urIPCOpenMemHandleExp" +returns: + - $X_RESULT_ERROR_INVALID_CONTEXT + - $X_RESULT_ERROR_INVALID_NULL_HANDLE: + - "`NULL == hContext`" + - $X_RESULT_ERROR_INVALID_NULL_POINTER: + - "`NULL == pMem`" + - $X_RESULT_ERROR_OUT_OF_HOST_MEMORY + - $X_RESULT_ERROR_OUT_OF_RESOURCES +--- #-------------------------------------------------------------------------- +type: function +desc: "Gets the data of an inter-process memory handle" +class: $xIPC +name: GetMemHandleDataExp +ordinal: "0" +params: + - type: $x_context_handle_t + name: hContext + desc: "[in] handle of the context object" + - type: $x_exp_ipc_mem_handle_t + name: hIPCMem + desc: "[in] the IPC memory handle" + - type: void** + name: ppIPCHandleData + desc: "[out][optional] a pointer to the IPC memory handle data" + - type: size_t* + name: pIPCMemHandleDataSizeRet + desc: "[out][optional] size of the resulting IPC memory handle data" +returns: + - $X_RESULT_ERROR_INVALID_CONTEXT + - $X_RESULT_ERROR_INVALID_NULL_HANDLE: + - "`NULL == hContext`" + - "`NULL == hIPCMem`" + - $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 a6237d93bf5ce..020374a3484bb 100644 --- a/unified-runtime/scripts/core/registry.yml +++ b/unified-runtime/scripts/core/registry.yml @@ -670,6 +670,21 @@ etors: - name: BINDLESS_IMAGES_SUPPORTS_IMPORTING_HANDLE_TYPE_EXP desc: Enumerator for $xBindlessImagesSupportsImportingHandleTypeExp value: '288' +- name: IPC_GET_MEM_HANDLE_EXP + desc: Enumerator for $xIPCGetMemHandleExp + value: '289' +- name: IPC_PUT_MEM_HANDLE_EXP + desc: Enumerator for $xIPCPutMemHandleExp + value: '290' +- name: IPC_OPEN_MEM_HANDLE_EXP + desc: Enumerator for $xIPCOpenMemHandleExp + value: '291' +- name: IPC_CLOSE_MEM_HANDLE_EXP + desc: Enumerator for $xIPCCloseMemHandleExp + value: '292' +- name: IPC_GET_MEM_HANDLE_DATA_EXP + desc: Enumerator for $xIPCGetMemHandleDataExp + value: '293' --- type: enum desc: Defines structure types diff --git a/unified-runtime/source/adapters/adapter.def.in b/unified-runtime/source/adapters/adapter.def.in index edccce444d2e0..aa3b55765ce3c 100644 --- a/unified-runtime/source/adapters/adapter.def.in +++ b/unified-runtime/source/adapters/adapter.def.in @@ -7,6 +7,7 @@ EXPORTS urGetEnqueueProcAddrTable urGetEnqueueExpProcAddrTable urGetEventProcAddrTable + urGetIPCExpProcAddrTable urGetKernelProcAddrTable urGetMemProcAddrTable urGetMemoryExportExpProcAddrTable diff --git a/unified-runtime/source/adapters/adapter.map.in b/unified-runtime/source/adapters/adapter.map.in index 54ff7d6b93f31..4be9b9b1348b7 100644 --- a/unified-runtime/source/adapters/adapter.map.in +++ b/unified-runtime/source/adapters/adapter.map.in @@ -7,6 +7,7 @@ urGetEnqueueProcAddrTable; urGetEnqueueExpProcAddrTable; urGetEventProcAddrTable; + urGetIPCExpProcAddrTable; urGetKernelProcAddrTable; urGetMemProcAddrTable; urGetMemoryExportExpProcAddrTable; diff --git a/unified-runtime/source/adapters/cuda/device.cpp b/unified-runtime/source/adapters/cuda/device.cpp index 03d9a13999f84..ddc4c7a115e26 100644 --- a/unified-runtime/source/adapters/cuda/device.cpp +++ b/unified-runtime/source/adapters/cuda/device.cpp @@ -1146,6 +1146,13 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_2D_BLOCK_ARRAY_CAPABILITIES_EXP: return ReturnValue( static_cast(0)); + case UR_DEVICE_INFO_IPC_MEMORY_SUPPORT_EXP: { + int IPCSupported = 0; + UR_CHECK_ERROR(cuDeviceGetAttribute(&IPCSupported, + CU_DEVICE_ATTRIBUTE_IPC_EVENT_SUPPORTED, + hDevice->get())); + return ReturnValue(static_cast(IPCSupported)); + } 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 b48252de6a816..0461b553a4fb0 100644 --- a/unified-runtime/source/adapters/cuda/memory.cpp +++ b/unified-runtime/source/adapters/cuda/memory.cpp @@ -15,6 +15,7 @@ #include "enqueue.hpp" #include "memory.hpp" #include "umf_helpers.hpp" +#include "usm.hpp" /// Creates a UR Memory object using a CUDA memory allocation. /// Can trigger a manual copy depending on the mode. @@ -589,3 +590,72 @@ CUsurfObject SurfaceMem::getSurface(const ur_device_handle_t Device) { } return SurfObjs[OuterMemStruct->getContext()->getDeviceIndex(Device)]; } + +UR_APIEXPORT ur_result_t UR_APICALL urIPCGetMemHandleExp( + ur_context_handle_t, void *pMem, ur_exp_ipc_mem_handle_t *phIPCMem) { + auto resHandle = std::make_unique(); + + umf_memory_pool_handle_t umfPool; + auto umfRet = umfPoolByPtr(pMem, &umfPool); + if (umfRet != UMF_RESULT_SUCCESS || !umfPool) + return UR_RESULT_ERROR_UNKNOWN; + + umfRet = umfGetIPCHandle(pMem, &resHandle->UMFHandle, &resHandle->HandleSize); + if (umfRet != UMF_RESULT_SUCCESS || !resHandle->UMFHandle || + resHandle->HandleSize == 0) + return UR_RESULT_ERROR_UNKNOWN; + + *phIPCMem = resHandle.release(); + return UR_RESULT_SUCCESS; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urIPCPutMemHandleExp(ur_context_handle_t, ur_exp_ipc_mem_handle_t hIPCMem) { + auto umfRet = umfPutIPCHandle(hIPCMem->UMFHandle); + if (umfRet != UMF_RESULT_SUCCESS) + return UR_RESULT_ERROR_UNKNOWN; + std::free(hIPCMem); + return UR_RESULT_SUCCESS; +} + +UR_APIEXPORT ur_result_t UR_APICALL urIPCOpenMemHandleExp( + ur_context_handle_t hContext, ur_device_handle_t hDevice, + void *pIPCMemHandleData, size_t ipcMemHandleDataSize, void **ppMem) { + umf_memory_pool_handle_t umfPool = hDevice->MemoryPoolDevice; + + size_t umfHandleSize = 0; + auto umfRet = umfPoolGetIPCHandleSize(umfPool, &umfHandleSize); + if (umfRet != UMF_RESULT_SUCCESS || umfHandleSize == 0) + return UR_RESULT_ERROR_UNKNOWN; + + if (umfHandleSize != ipcMemHandleDataSize) + return UR_RESULT_ERROR_INVALID_VALUE; + + umf_ipc_handler_handle_t umfIPCHandler; + umfRet = umfPoolGetIPCHandler(umfPool, &umfIPCHandler); + if (umfRet != UMF_RESULT_SUCCESS || !umfIPCHandler) + return UR_RESULT_ERROR_UNKNOWN; + + umfRet = umfOpenIPCHandle( + umfIPCHandler, reinterpret_cast(pIPCMemHandleData), + ppMem); + return umfRet == UMF_RESULT_SUCCESS ? UR_RESULT_SUCCESS + : UR_RESULT_ERROR_UNKNOWN; +} + +UR_APIEXPORT ur_result_t UR_APICALL urIPCCloseMemHandleExp(ur_context_handle_t, + void *pMem) { + auto umfRet = umfCloseIPCHandle(pMem); + return umfRet == UMF_RESULT_SUCCESS ? UR_RESULT_SUCCESS + : UR_RESULT_ERROR_UNKNOWN; +} + +UR_APIEXPORT ur_result_t UR_APICALL urIPCGetMemHandleDataExp( + ur_context_handle_t, ur_exp_ipc_mem_handle_t hIPCMem, + void **ppIPCHandleData, size_t *pIPCMemHandleDataSizeRet) { + if (ppIPCHandleData) + *ppIPCHandleData = hIPCMem->UMFHandle; + if (pIPCMemHandleDataSizeRet) + *pIPCMemHandleDataSizeRet = hIPCMem->HandleSize; + return UR_RESULT_SUCCESS; +} diff --git a/unified-runtime/source/adapters/cuda/memory.hpp b/unified-runtime/source/adapters/cuda/memory.hpp index 6bb817f1efee8..9d3f19c4b9c58 100644 --- a/unified-runtime/source/adapters/cuda/memory.hpp +++ b/unified-runtime/source/adapters/cuda/memory.hpp @@ -19,6 +19,7 @@ #include "common/ur_ref_count.hpp" #include "context.hpp" #include "queue.hpp" +#include ur_result_t allocateMemObjOnDeviceIfNeeded(ur_mem_handle_t, const ur_device_handle_t); @@ -439,3 +440,8 @@ struct ur_mem_handle_t_ : ur::cuda::handle_base { } } }; + +struct ur_exp_ipc_mem_handle_t_ { + umf_ipc_handle_t UMFHandle = nullptr; + size_t HandleSize = 0; +}; diff --git a/unified-runtime/source/adapters/cuda/ur_interface_loader.cpp b/unified-runtime/source/adapters/cuda/ur_interface_loader.cpp index 8430df0ab0678..2eb42d8bc0dea 100644 --- a/unified-runtime/source/adapters/cuda/ur_interface_loader.cpp +++ b/unified-runtime/source/adapters/cuda/ur_interface_loader.cpp @@ -458,6 +458,22 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetEnqueueExpProcAddrTable( return UR_RESULT_SUCCESS; } +UR_DLLEXPORT ur_result_t UR_APICALL urGetIPCExpProcAddrTable( + ur_api_version_t version, ur_ipc_exp_dditable_t *pDdiTable) { + auto result = validateProcInputs(version, pDdiTable); + if (UR_RESULT_SUCCESS != result) { + return result; + } + + pDdiTable->pfnGetMemHandleExp = urIPCGetMemHandleExp; + pDdiTable->pfnPutMemHandleExp = urIPCPutMemHandleExp; + pDdiTable->pfnOpenMemHandleExp = urIPCOpenMemHandleExp; + pDdiTable->pfnCloseMemHandleExp = urIPCCloseMemHandleExp; + pDdiTable->pfnGetMemHandleDataExp = urIPCGetMemHandleDataExp; + + return UR_RESULT_SUCCESS; +} + UR_DLLEXPORT ur_result_t UR_APICALL urGetProgramExpProcAddrTable( ur_api_version_t version, ur_program_exp_dditable_t *pDdiTable) { auto result = validateProcInputs(version, pDdiTable); @@ -480,6 +496,7 @@ UR_DLLEXPORT ur_result_t UR_APICALL urAllAddrTable(ur_api_version_t version, urGetContextProcAddrTable(version, &pDdiTable->Context); urGetEnqueueProcAddrTable(version, &pDdiTable->Enqueue); urGetEnqueueExpProcAddrTable(version, &pDdiTable->EnqueueExp); + urGetIPCExpProcAddrTable(version, &pDdiTable->IPCExp); urGetEventProcAddrTable(version, &pDdiTable->Event); urGetKernelProcAddrTable(version, &pDdiTable->Kernel); urGetMemProcAddrTable(version, &pDdiTable->Mem); diff --git a/unified-runtime/source/adapters/hip/device.cpp b/unified-runtime/source/adapters/hip/device.cpp index c48033ec88826..f71d86c3f4a5e 100644 --- a/unified-runtime/source/adapters/hip/device.cpp +++ b/unified-runtime/source/adapters/hip/device.cpp @@ -999,6 +999,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_2D_BLOCK_ARRAY_CAPABILITIES_EXP: return ReturnValue( static_cast(0)); + case UR_DEVICE_INFO_IPC_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 7995b15d36f84..f052fd2efc785 100644 --- a/unified-runtime/source/adapters/hip/memory.cpp +++ b/unified-runtime/source/adapters/hip/memory.cpp @@ -640,3 +640,30 @@ hipSurfaceObject_t SurfaceMem::getSurface(const ur_device_handle_t Device) { } return SurfObjs[OuterMemStruct->getContext()->getDeviceIndex(Device)]; } + +UR_APIEXPORT ur_result_t UR_APICALL +urIPCGetMemHandleExp(ur_context_handle_t, void *, ur_exp_ipc_mem_handle_t *) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urIPCPutMemHandleExp(ur_context_handle_t, ur_exp_ipc_mem_handle_t) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +UR_APIEXPORT ur_result_t UR_APICALL urIPCOpenMemHandleExp(ur_context_handle_t, + ur_device_handle_t, + void *, size_t, + void **) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +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 urIPCGetMemHandleDataExp( + ur_context_handle_t, ur_exp_ipc_mem_handle_t, void **, size_t *) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} diff --git a/unified-runtime/source/adapters/hip/memory.hpp b/unified-runtime/source/adapters/hip/memory.hpp index 239e3949c740c..d825d153b04ce 100644 --- a/unified-runtime/source/adapters/hip/memory.hpp +++ b/unified-runtime/source/adapters/hip/memory.hpp @@ -433,3 +433,6 @@ struct ur_mem_handle_t_ : ur::hip::handle_base { } } }; + +// IPC is currently not supported in the HIP adaptor. +struct ur_exp_ipc_mem_handle_t_ {}; diff --git a/unified-runtime/source/adapters/hip/ur_interface_loader.cpp b/unified-runtime/source/adapters/hip/ur_interface_loader.cpp index dfb4382cad828..9c6cf548b4198 100644 --- a/unified-runtime/source/adapters/hip/ur_interface_loader.cpp +++ b/unified-runtime/source/adapters/hip/ur_interface_loader.cpp @@ -451,6 +451,22 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetEnqueueExpProcAddrTable( return UR_RESULT_SUCCESS; } +UR_DLLEXPORT ur_result_t UR_APICALL urGetIPCExpProcAddrTable( + ur_api_version_t version, ur_ipc_exp_dditable_t *pDdiTable) { + auto result = validateProcInputs(version, pDdiTable); + if (UR_RESULT_SUCCESS != result) { + return result; + } + + pDdiTable->pfnGetMemHandleExp = urIPCGetMemHandleExp; + pDdiTable->pfnPutMemHandleExp = urIPCPutMemHandleExp; + pDdiTable->pfnOpenMemHandleExp = urIPCOpenMemHandleExp; + pDdiTable->pfnCloseMemHandleExp = urIPCCloseMemHandleExp; + pDdiTable->pfnGetMemHandleDataExp = urIPCGetMemHandleDataExp; + + return UR_RESULT_SUCCESS; +} + UR_DLLEXPORT ur_result_t UR_APICALL urGetProgramExpProcAddrTable( ur_api_version_t version, ur_program_exp_dditable_t *pDdiTable) { auto result = validateProcInputs(version, pDdiTable); @@ -473,6 +489,7 @@ UR_DLLEXPORT ur_result_t UR_APICALL urAllAddrTable(ur_api_version_t version, urGetContextProcAddrTable(version, &pDdiTable->Context); urGetEnqueueProcAddrTable(version, &pDdiTable->Enqueue); urGetEnqueueExpProcAddrTable(version, &pDdiTable->EnqueueExp); + urGetIPCExpProcAddrTable(version, &pDdiTable->IPCExp); urGetEventProcAddrTable(version, &pDdiTable->Event); urGetKernelProcAddrTable(version, &pDdiTable->Kernel); urGetMemProcAddrTable(version, &pDdiTable->Mem); diff --git a/unified-runtime/source/adapters/level_zero/device.cpp b/unified-runtime/source/adapters/level_zero/device.cpp index a00d816d5ed61..37d488bd34239 100644 --- a/unified-runtime/source/adapters/level_zero/device.cpp +++ b/unified-runtime/source/adapters/level_zero/device.cpp @@ -1276,6 +1276,8 @@ ur_result_t urDeviceGetInfo( return UR_RESULT_ERROR_UNSUPPORTED_ENUMERATION; #endif } + case UR_DEVICE_INFO_IPC_MEMORY_SUPPORT_EXP: + return ReturnValue(true); case UR_DEVICE_INFO_ASYNC_BARRIER: return ReturnValue(false); case UR_DEVICE_INFO_HOST_PIPE_READ_WRITE_SUPPORT: diff --git a/unified-runtime/source/adapters/level_zero/memory.cpp b/unified-runtime/source/adapters/level_zero/memory.cpp index 107fcc2d1c2f5..6c977dace0113 100644 --- a/unified-runtime/source/adapters/level_zero/memory.cpp +++ b/unified-runtime/source/adapters/level_zero/memory.cpp @@ -1952,6 +1952,80 @@ ur_result_t urEnqueueWriteHostPipe( return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } +ur_result_t urIPCGetMemHandleExp(ur_context_handle_t, void *pMem, + ur_exp_ipc_mem_handle_t *phIPCMem) { + auto resHandle = std::make_unique(); + + umf_memory_pool_handle_t umfPool; + auto umfRet = umfPoolByPtr(pMem, &umfPool); + if (umfRet != UMF_RESULT_SUCCESS || !umfPool) + return UR_RESULT_ERROR_UNKNOWN; + + umfRet = umfGetIPCHandle(pMem, &resHandle->UMFHandle, &resHandle->HandleSize); + if (umfRet != UMF_RESULT_SUCCESS || !resHandle->UMFHandle || + resHandle->HandleSize == 0) + return UR_RESULT_ERROR_UNKNOWN; + + *phIPCMem = resHandle.release(); + return UR_RESULT_SUCCESS; +} + +ur_result_t urIPCPutMemHandleExp(ur_context_handle_t, + ur_exp_ipc_mem_handle_t hIPCMem) { + auto umfRet = umfPutIPCHandle(hIPCMem->UMFHandle); + if (umfRet != UMF_RESULT_SUCCESS) + return UR_RESULT_ERROR_UNKNOWN; + std::free(hIPCMem); + return UR_RESULT_SUCCESS; +} + +ur_result_t urIPCOpenMemHandleExp(ur_context_handle_t hContext, + ur_device_handle_t hDevice, + void *pIPCMemHandleData, + size_t ipcMemHandleDataSize, void **ppMem) { + auto *pool = hContext->DefaultPool.getPool(usm::pool_descriptor{ + &hContext->DefaultPool, hContext, hDevice, UR_USM_TYPE_DEVICE, false}); + if (!pool) + return UR_RESULT_ERROR_INVALID_CONTEXT; + umf_memory_pool_handle_t umfPool = pool->UmfPool.get(); + + size_t umfHandleSize = 0; + auto umfRet = umfPoolGetIPCHandleSize(umfPool, &umfHandleSize); + if (umfRet != UMF_RESULT_SUCCESS || umfHandleSize == 0) + return UR_RESULT_ERROR_UNKNOWN; + + if (umfHandleSize != ipcMemHandleDataSize) + return UR_RESULT_ERROR_INVALID_VALUE; + + umf_ipc_handler_handle_t umfIPCHandler; + umfRet = umfPoolGetIPCHandler(umfPool, &umfIPCHandler); + if (umfRet != UMF_RESULT_SUCCESS || !umfIPCHandler) + return UR_RESULT_ERROR_UNKNOWN; + + umfRet = umfOpenIPCHandle( + umfIPCHandler, reinterpret_cast(pIPCMemHandleData), + ppMem); + return umfRet == UMF_RESULT_SUCCESS ? UR_RESULT_SUCCESS + : UR_RESULT_ERROR_UNKNOWN; +} + +ur_result_t urIPCCloseMemHandleExp(ur_context_handle_t, void *pMem) { + auto umfRet = umfCloseIPCHandle(pMem); + return umfRet == UMF_RESULT_SUCCESS ? UR_RESULT_SUCCESS + : UR_RESULT_ERROR_UNKNOWN; +} + +ur_result_t urIPCGetMemHandleDataExp(ur_context_handle_t, + ur_exp_ipc_mem_handle_t hIPCMem, + void **ppIPCHandleData, + size_t *pIPCMemHandleDataSizeRet) { + if (ppIPCHandleData) + *ppIPCHandleData = hIPCMem->UMFHandle; + if (pIPCMemHandleDataSizeRet) + *pIPCMemHandleDataSizeRet = hIPCMem->HandleSize; + return UR_RESULT_SUCCESS; +} + } // namespace ur::level_zero // If indirect access tracking is enabled then performs reference counting, diff --git a/unified-runtime/source/adapters/level_zero/memory.hpp b/unified-runtime/source/adapters/level_zero/memory.hpp index f58f189b21c77..cc3917f0e6b43 100644 --- a/unified-runtime/source/adapters/level_zero/memory.hpp +++ b/unified-runtime/source/adapters/level_zero/memory.hpp @@ -25,6 +25,7 @@ #include "program.hpp" #include "queue.hpp" #include "sampler.hpp" +#include struct ur_device_handle_t_; @@ -243,6 +244,11 @@ struct ur_image final : ur_mem_handle_t_ { ze_image_handle_t ZeImage; }; +struct ur_exp_ipc_mem_handle_t_ { + umf_ipc_handle_t UMFHandle = nullptr; + size_t HandleSize = 0; +}; + template ur_result_t createUrMemFromZeImage(ur_context_handle_t Context, ze_image_handle_t ZeImage, 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 13d7274e7aebf..fe2d9d49e6024 100644 --- a/unified-runtime/source/adapters/level_zero/ur_interface_loader.cpp +++ b/unified-runtime/source/adapters/level_zero/ur_interface_loader.cpp @@ -257,6 +257,22 @@ UR_APIEXPORT ur_result_t UR_APICALL urGetEventProcAddrTable( return result; } +UR_APIEXPORT ur_result_t UR_APICALL urGetIPCExpProcAddrTable( + ur_api_version_t version, ur_ipc_exp_dditable_t *pDdiTable) { + auto result = validateProcInputs(version, pDdiTable); + if (UR_RESULT_SUCCESS != result) { + return result; + } + + pDdiTable->pfnGetMemHandleExp = ur::level_zero::urIPCGetMemHandleExp; + pDdiTable->pfnPutMemHandleExp = ur::level_zero::urIPCPutMemHandleExp; + pDdiTable->pfnOpenMemHandleExp = ur::level_zero::urIPCOpenMemHandleExp; + pDdiTable->pfnCloseMemHandleExp = ur::level_zero::urIPCCloseMemHandleExp; + pDdiTable->pfnGetMemHandleDataExp = ur::level_zero::urIPCGetMemHandleDataExp; + + return result; +} + UR_APIEXPORT ur_result_t UR_APICALL urGetKernelProcAddrTable( ur_api_version_t version, ur_kernel_dditable_t *pDdiTable) { auto result = validateProcInputs(version, pDdiTable); @@ -595,6 +611,10 @@ ur_result_t populateDdiTable(ur_dditable_t *ddi) { NAMESPACE_::urGetEventProcAddrTable(UR_API_VERSION_CURRENT, &ddi->Event); if (result != UR_RESULT_SUCCESS) return result; + result = NAMESPACE_::urGetIPCExpProcAddrTable(UR_API_VERSION_CURRENT, + &ddi->IPCExp); + if (result != UR_RESULT_SUCCESS) + return result; result = NAMESPACE_::urGetKernelProcAddrTable(UR_API_VERSION_CURRENT, &ddi->Kernel); if (result != UR_RESULT_SUCCESS) 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 df8e93c1f768a..0c7f6a9694be0 100644 --- a/unified-runtime/source/adapters/level_zero/ur_interface_loader.hpp +++ b/unified-runtime/source/adapters/level_zero/ur_interface_loader.hpp @@ -768,6 +768,19 @@ urCommandBufferGetNativeHandleExp(ur_exp_command_buffer_handle_t hCommandBuffer, ur_result_t urEnqueueTimestampRecordingExp( ur_queue_handle_t hQueue, bool blocking, uint32_t numEventsInWaitList, const ur_event_handle_t *phEventWaitList, ur_event_handle_t *phEvent); +ur_result_t urIPCGetMemHandleExp(ur_context_handle_t hContext, void *pMem, + ur_exp_ipc_mem_handle_t *phIPCMem); +ur_result_t urIPCPutMemHandleExp(ur_context_handle_t hContext, + ur_exp_ipc_mem_handle_t hIPCMem); +ur_result_t urIPCOpenMemHandleExp(ur_context_handle_t hContext, + ur_device_handle_t hDevice, + void *ipcMemHandleData, + size_t ipcMemHandleDataSize, void **ppMem); +ur_result_t urIPCCloseMemHandleExp(ur_context_handle_t hContext, void *pMem); +ur_result_t urIPCGetMemHandleDataExp(ur_context_handle_t hContext, + ur_exp_ipc_mem_handle_t hIPCMem, + void **ppIPCHandleData, + size_t *pIPCMemHandleDataSizeRet); 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/usm.hpp b/unified-runtime/source/adapters/level_zero/usm.hpp index aa7558926b5ed..d3e2e344f149f 100644 --- a/unified-runtime/source/adapters/level_zero/usm.hpp +++ b/unified-runtime/source/adapters/level_zero/usm.hpp @@ -84,12 +84,12 @@ struct ur_usm_pool_handle_t_ : ur_object { size_t getPeakReservedSize(); size_t getTotalUsedSize(); size_t getPeakUsedSize(); + UsmPool *getPool(const usm::pool_descriptor &Desc); ur_context_handle_t Context; ur::RefCount RefCount; private: - UsmPool *getPool(const usm::pool_descriptor &Desc); usm::pool_manager PoolManager; AllocationStats AllocStats; }; diff --git a/unified-runtime/source/adapters/level_zero/v2/memory.cpp b/unified-runtime/source/adapters/level_zero/v2/memory.cpp index 1b6855e630994..1794cfe0cffc3 100644 --- a/unified-runtime/source/adapters/level_zero/v2/memory.cpp +++ b/unified-runtime/source/adapters/level_zero/v2/memory.cpp @@ -775,4 +775,79 @@ ur_result_t urMemImageGetInfo(ur_mem_handle_t /*hMemory*/, return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } +ur_result_t urIPCGetMemHandleExp(ur_context_handle_t, void *pMem, + ur_exp_ipc_mem_handle_t *phIPCMem) { + auto resHandle = std::make_unique(); + + umf_memory_pool_handle_t umfPool; + auto umfRet = umfPoolByPtr(pMem, &umfPool); + if (umfRet != UMF_RESULT_SUCCESS || !umfPool) + return UR_RESULT_ERROR_UNKNOWN; + + umfRet = umfGetIPCHandle(pMem, &resHandle->UMFHandle, &resHandle->HandleSize); + if (umfRet != UMF_RESULT_SUCCESS || !resHandle->UMFHandle || + resHandle->HandleSize == 0) + return UR_RESULT_ERROR_UNKNOWN; + + *phIPCMem = resHandle.release(); + return UR_RESULT_SUCCESS; +} + +ur_result_t urIPCPutMemHandleExp(ur_context_handle_t, + ur_exp_ipc_mem_handle_t hIPCMem) { + auto umfRet = umfPutIPCHandle(hIPCMem->UMFHandle); + if (umfRet != UMF_RESULT_SUCCESS) + return UR_RESULT_ERROR_UNKNOWN; + std::free(hIPCMem); + return UR_RESULT_SUCCESS; +} + +ur_result_t urIPCOpenMemHandleExp(ur_context_handle_t hContext, + ur_device_handle_t hDevice, + void *pIPCMemHandleData, + size_t ipcMemHandleDataSize, void **ppMem) { + auto *pool = hContext->getDefaultUSMPool()->getPool( + usm::pool_descriptor{hContext->getDefaultUSMPool(), hContext, hDevice, + UR_USM_TYPE_DEVICE, false}); + if (!pool) + return UR_RESULT_ERROR_INVALID_CONTEXT; + umf_memory_pool_handle_t umfPool = pool->umfPool.get(); + + size_t umfHandleSize = 0; + auto umfRet = umfPoolGetIPCHandleSize(umfPool, &umfHandleSize); + if (umfRet != UMF_RESULT_SUCCESS || umfHandleSize == 0) + return UR_RESULT_ERROR_UNKNOWN; + + if (umfHandleSize != ipcMemHandleDataSize) + return UR_RESULT_ERROR_INVALID_VALUE; + + umf_ipc_handler_handle_t umfIPCHandler; + umfRet = umfPoolGetIPCHandler(umfPool, &umfIPCHandler); + if (umfRet != UMF_RESULT_SUCCESS || !umfIPCHandler) + return UR_RESULT_ERROR_UNKNOWN; + + umfRet = umfOpenIPCHandle( + umfIPCHandler, reinterpret_cast(pIPCMemHandleData), + ppMem); + return umfRet == UMF_RESULT_SUCCESS ? UR_RESULT_SUCCESS + : UR_RESULT_ERROR_UNKNOWN; +} + +ur_result_t urIPCCloseMemHandleExp(ur_context_handle_t, void *pMem) { + auto umfRet = umfCloseIPCHandle(pMem); + return umfRet == UMF_RESULT_SUCCESS ? UR_RESULT_SUCCESS + : UR_RESULT_ERROR_UNKNOWN; +} + +ur_result_t urIPCGetMemHandleDataExp(ur_context_handle_t, + ur_exp_ipc_mem_handle_t hIPCMem, + void **ppIPCHandleData, + size_t *pIPCMemHandleDataSizeRet) { + if (ppIPCHandleData) + *ppIPCHandleData = hIPCMem->UMFHandle; + if (pIPCMemHandleDataSizeRet) + *pIPCMemHandleDataSizeRet = hIPCMem->HandleSize; + return UR_RESULT_SUCCESS; +} + } // namespace ur::level_zero diff --git a/unified-runtime/source/adapters/level_zero/v2/memory.hpp b/unified-runtime/source/adapters/level_zero/v2/memory.hpp index 61b0a00f4043b..c61a171c511c4 100644 --- a/unified-runtime/source/adapters/level_zero/v2/memory.hpp +++ b/unified-runtime/source/adapters/level_zero/v2/memory.hpp @@ -20,6 +20,7 @@ #include "command_list_manager.hpp" #include "common.hpp" #include "common/ur_ref_count.hpp" +#include using usm_unique_ptr_t = std::unique_ptr>; @@ -292,3 +293,8 @@ struct ur_mem_handle_t_ : ur::handle_base { ur_mem_sub_buffer_t, ur_mem_image_t> mem; }; + +struct ur_exp_ipc_mem_handle_t_ { + umf_ipc_handle_t UMFHandle = nullptr; + size_t HandleSize = 0; +}; diff --git a/unified-runtime/source/adapters/level_zero/v2/usm.hpp b/unified-runtime/source/adapters/level_zero/v2/usm.hpp index 825ecb5fcd8e3..ace4b7411e617 100644 --- a/unified-runtime/source/adapters/level_zero/v2/usm.hpp +++ b/unified-runtime/source/adapters/level_zero/v2/usm.hpp @@ -81,14 +81,14 @@ struct ur_usm_pool_handle_t_ : ur_object { size_t getTotalUsedSize(); size_t getPeakUsedSize(); + UsmPool *getPool(const usm::pool_descriptor &desc); + ur::RefCount RefCount; private: ur_context_handle_t hContext; usm::pool_manager poolManager; AllocationStats allocStats; - - UsmPool *getPool(const usm::pool_descriptor &desc); }; struct UsmPool { diff --git a/unified-runtime/source/adapters/mock/ur_mock.cpp b/unified-runtime/source/adapters/mock/ur_mock.cpp index 142a706146a9b..bd2112793a3a1 100644 --- a/unified-runtime/source/adapters/mock/ur_mock.cpp +++ b/unified-runtime/source/adapters/mock/ur_mock.cpp @@ -103,6 +103,7 @@ context_t::context_t() { urGetContextProcAddrTable(version, &urDdiTable.Context); urGetEnqueueProcAddrTable(version, &urDdiTable.Enqueue); urGetEnqueueExpProcAddrTable(version, &urDdiTable.EnqueueExp); + urGetIPCExpProcAddrTable(version, &urDdiTable.IPCExp); urGetEventProcAddrTable(version, &urDdiTable.Event); urGetKernelProcAddrTable(version, &urDdiTable.Kernel); urGetMemProcAddrTable(version, &urDdiTable.Mem); diff --git a/unified-runtime/source/adapters/mock/ur_mockddi.cpp b/unified-runtime/source/adapters/mock/ur_mockddi.cpp index 7956f048db92e..80fd8024595e0 100644 --- a/unified-runtime/source/adapters/mock/ur_mockddi.cpp +++ b/unified-runtime/source/adapters/mock/ur_mockddi.cpp @@ -11304,6 +11304,241 @@ __urdlllocal ur_result_t UR_APICALL urEnqueueTimestampRecordingExp( return exceptionToResult(std::current_exception()); } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCGetMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCGetMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] pointer to device USM memory + void *pMem, + /// [out][alloc] pointer to the resulting IPC memory handle + ur_exp_ipc_mem_handle_t *phIPCMem) try { + ur_result_t result = UR_RESULT_SUCCESS; + + ur_ipc_get_mem_handle_exp_params_t params = {&hContext, &pMem, &phIPCMem}; + + auto beforeCallback = reinterpret_cast( + mock::getCallbacks().get_before_callback("urIPCGetMemHandleExp")); + if (beforeCallback) { + result = beforeCallback(¶ms); + if (result != UR_RESULT_SUCCESS) { + return result; + } + } + + auto replaceCallback = reinterpret_cast( + mock::getCallbacks().get_replace_callback("urIPCGetMemHandleExp")); + 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("urIPCGetMemHandleExp")); + if (afterCallback) { + return afterCallback(¶ms); + } + + return result; +} catch (...) { + return exceptionToResult(std::current_exception()); +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCPutMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCPutMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] the IPC memory handle + ur_exp_ipc_mem_handle_t hIPCMem) try { + ur_result_t result = UR_RESULT_SUCCESS; + + ur_ipc_put_mem_handle_exp_params_t params = {&hContext, &hIPCMem}; + + auto beforeCallback = reinterpret_cast( + mock::getCallbacks().get_before_callback("urIPCPutMemHandleExp")); + if (beforeCallback) { + result = beforeCallback(¶ms); + if (result != UR_RESULT_SUCCESS) { + return result; + } + } + + auto replaceCallback = reinterpret_cast( + mock::getCallbacks().get_replace_callback("urIPCPutMemHandleExp")); + 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("urIPCPutMemHandleExp")); + if (afterCallback) { + return afterCallback(¶ms); + } + + return result; +} catch (...) { + return exceptionToResult(std::current_exception()); +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCOpenMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCOpenMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object the corresponding USM device memory + /// was allocated on + ur_device_handle_t hDevice, + /// [in] the IPC memory handle data + void *ipcMemHandleData, + /// [in] size of the IPC memory handle data + size_t ipcMemHandleDataSize, + /// [out] pointer to a pointer to device USM memory + void **ppMem) try { + ur_result_t result = UR_RESULT_SUCCESS; + + ur_ipc_open_mem_handle_exp_params_t params = { + &hContext, &hDevice, &ipcMemHandleData, &ipcMemHandleDataSize, &ppMem}; + + auto beforeCallback = reinterpret_cast( + mock::getCallbacks().get_before_callback("urIPCOpenMemHandleExp")); + if (beforeCallback) { + result = beforeCallback(¶ms); + if (result != UR_RESULT_SUCCESS) { + return result; + } + } + + auto replaceCallback = reinterpret_cast( + mock::getCallbacks().get_replace_callback("urIPCOpenMemHandleExp")); + 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("urIPCOpenMemHandleExp")); + if (afterCallback) { + return afterCallback(¶ms); + } + + return result; +} catch (...) { + return exceptionToResult(std::current_exception()); +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCCloseMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCCloseMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] pointer to device USM memory opened through urIPCOpenMemHandleExp + void *pMem) try { + ur_result_t result = UR_RESULT_SUCCESS; + + ur_ipc_close_mem_handle_exp_params_t params = {&hContext, &pMem}; + + auto beforeCallback = reinterpret_cast( + mock::getCallbacks().get_before_callback("urIPCCloseMemHandleExp")); + if (beforeCallback) { + result = beforeCallback(¶ms); + if (result != UR_RESULT_SUCCESS) { + return result; + } + } + + auto replaceCallback = reinterpret_cast( + mock::getCallbacks().get_replace_callback("urIPCCloseMemHandleExp")); + 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("urIPCCloseMemHandleExp")); + if (afterCallback) { + return afterCallback(¶ms); + } + + return result; +} catch (...) { + return exceptionToResult(std::current_exception()); +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCGetMemHandleDataExp +__urdlllocal ur_result_t UR_APICALL urIPCGetMemHandleDataExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] the IPC memory handle + ur_exp_ipc_mem_handle_t hIPCMem, + /// [out][optional] a pointer to the IPC memory handle data + void **ppIPCHandleData, + /// [out][optional] size of the resulting IPC memory handle data + size_t *pIPCMemHandleDataSizeRet) try { + ur_result_t result = UR_RESULT_SUCCESS; + + ur_ipc_get_mem_handle_data_exp_params_t params = { + &hContext, &hIPCMem, &ppIPCHandleData, &pIPCMemHandleDataSizeRet}; + + auto beforeCallback = reinterpret_cast( + mock::getCallbacks().get_before_callback("urIPCGetMemHandleDataExp")); + if (beforeCallback) { + result = beforeCallback(¶ms); + if (result != UR_RESULT_SUCCESS) { + return result; + } + } + + auto replaceCallback = reinterpret_cast( + mock::getCallbacks().get_replace_callback("urIPCGetMemHandleDataExp")); + 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("urIPCGetMemHandleDataExp")); + 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( @@ -12481,6 +12716,42 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetEventProcAddrTable( return exceptionToResult(std::current_exception()); } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Exported function for filling application's IPCExp table +/// with current process' addresses +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// - ::UR_RESULT_ERROR_UNSUPPORTED_VERSION +UR_DLLEXPORT ur_result_t UR_APICALL urGetIPCExpProcAddrTable( + /// [in] API version requested + ur_api_version_t version, + /// [in,out] pointer to table of DDI function pointers + ur_ipc_exp_dditable_t *pDdiTable) try { + if (nullptr == pDdiTable) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + + if (driver::d_context.version < version) + return UR_RESULT_ERROR_UNSUPPORTED_VERSION; + + ur_result_t result = UR_RESULT_SUCCESS; + + pDdiTable->pfnGetMemHandleExp = driver::urIPCGetMemHandleExp; + + pDdiTable->pfnPutMemHandleExp = driver::urIPCPutMemHandleExp; + + pDdiTable->pfnOpenMemHandleExp = driver::urIPCOpenMemHandleExp; + + pDdiTable->pfnCloseMemHandleExp = driver::urIPCCloseMemHandleExp; + + pDdiTable->pfnGetMemHandleDataExp = driver::urIPCGetMemHandleDataExp; + + return result; +} catch (...) { + return exceptionToResult(std::current_exception()); +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Exported function for filling application's Kernel table /// with current process' addresses diff --git a/unified-runtime/source/adapters/native_cpu/device.cpp b/unified-runtime/source/adapters/native_cpu/device.cpp index 369b4cd7ed013..4fed1565f8cfd 100644 --- a/unified-runtime/source/adapters/native_cpu/device.cpp +++ b/unified-runtime/source/adapters/native_cpu/device.cpp @@ -376,6 +376,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, case UR_DEVICE_INFO_2D_BLOCK_ARRAY_CAPABILITIES_EXP: return ReturnValue( static_cast(0)); + case UR_DEVICE_INFO_IPC_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 67eb95f3d9523..f856b19843c78 100644 --- a/unified-runtime/source/adapters/native_cpu/memory.cpp +++ b/unified-runtime/source/adapters/native_cpu/memory.cpp @@ -138,3 +138,30 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemImageGetInfo( DIE_NO_IMPLEMENTATION; } + +UR_APIEXPORT ur_result_t UR_APICALL +urIPCGetMemHandleExp(ur_context_handle_t, void *, ur_exp_ipc_mem_handle_t *) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urIPCPutMemHandleExp(ur_context_handle_t, ur_exp_ipc_mem_handle_t) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +UR_APIEXPORT ur_result_t UR_APICALL urIPCOpenMemHandleExp(ur_context_handle_t, + ur_device_handle_t, + void *, size_t, + void **) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +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 urIPCGetMemHandleDataExp( + ur_context_handle_t, ur_exp_ipc_mem_handle_t, void **, size_t *) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} diff --git a/unified-runtime/source/adapters/native_cpu/memory.hpp b/unified-runtime/source/adapters/native_cpu/memory.hpp index ca6e3e77f5e87..6285287e2457b 100644 --- a/unified-runtime/source/adapters/native_cpu/memory.hpp +++ b/unified-runtime/source/adapters/native_cpu/memory.hpp @@ -69,3 +69,6 @@ struct ur_buffer final : ur_mem_handle_t_ { size_t Origin; // only valid if Parent != nullptr } SubBuffer; }; + +// IPC is currently not supported in the native CPU adaptor. +struct ur_exp_ipc_mem_handle_t_ {}; 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 3f6fe061b4917..f0c36bb6dadc2 100644 --- a/unified-runtime/source/adapters/native_cpu/ur_interface_loader.cpp +++ b/unified-runtime/source/adapters/native_cpu/ur_interface_loader.cpp @@ -435,6 +435,22 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetEnqueueExpProcAddrTable( return UR_RESULT_SUCCESS; } +UR_DLLEXPORT ur_result_t UR_APICALL urGetIPCExpProcAddrTable( + ur_api_version_t version, ur_ipc_exp_dditable_t *pDdiTable) { + auto result = validateProcInputs(version, pDdiTable); + if (UR_RESULT_SUCCESS != result) { + return result; + } + + pDdiTable->pfnGetMemHandleExp = urIPCGetMemHandleExp; + pDdiTable->pfnPutMemHandleExp = urIPCPutMemHandleExp; + pDdiTable->pfnOpenMemHandleExp = urIPCOpenMemHandleExp; + pDdiTable->pfnCloseMemHandleExp = urIPCCloseMemHandleExp; + pDdiTable->pfnGetMemHandleDataExp = urIPCGetMemHandleDataExp; + + return UR_RESULT_SUCCESS; +} + UR_DLLEXPORT ur_result_t UR_APICALL urGetProgramExpProcAddrTable( ur_api_version_t version, ur_program_exp_dditable_t *pDdiTable) { auto result = validateProcInputs(version, pDdiTable); @@ -457,6 +473,7 @@ UR_DLLEXPORT ur_result_t UR_APICALL urAllAddrTable(ur_api_version_t version, urGetContextProcAddrTable(version, &pDdiTable->Context); urGetEnqueueProcAddrTable(version, &pDdiTable->Enqueue); urGetEnqueueExpProcAddrTable(version, &pDdiTable->EnqueueExp); + urGetIPCExpProcAddrTable(version, &pDdiTable->IPCExp); urGetEventProcAddrTable(version, &pDdiTable->Event); urGetKernelProcAddrTable(version, &pDdiTable->Kernel); urGetMemProcAddrTable(version, &pDdiTable->Mem); diff --git a/unified-runtime/source/adapters/offload/memory.cpp b/unified-runtime/source/adapters/offload/memory.cpp index e27a032a61451..fdc9a1a1d4a65 100644 --- a/unified-runtime/source/adapters/offload/memory.cpp +++ b/unified-runtime/source/adapters/offload/memory.cpp @@ -142,3 +142,30 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemBufferPartition( return urMemRetain(hBuffer); } + +UR_APIEXPORT ur_result_t UR_APICALL +urIPCGetMemHandleExp(ur_context_handle_t, void *, ur_exp_ipc_mem_handle_t *) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urIPCPutMemHandleExp(ur_context_handle_t, ur_exp_ipc_mem_handle_t) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +UR_APIEXPORT ur_result_t UR_APICALL urIPCOpenMemHandleExp(ur_context_handle_t, + ur_device_handle_t, + void *, size_t, + void **) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +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 urIPCGetMemHandleDataExp( + ur_context_handle_t, ur_exp_ipc_mem_handle_t, void **, size_t *) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} diff --git a/unified-runtime/source/adapters/offload/memory.hpp b/unified-runtime/source/adapters/offload/memory.hpp index 59b62ea12961a..a4742d7b288d3 100644 --- a/unified-runtime/source/adapters/offload/memory.hpp +++ b/unified-runtime/source/adapters/offload/memory.hpp @@ -117,3 +117,6 @@ struct ur_mem_handle_t_ : RefCounted { return nullptr; } }; + +// IPC is currently not supported in the offload adaptor. +struct ur_exp_ipc_mem_handle_t_ {}; diff --git a/unified-runtime/source/adapters/offload/ur_interface_loader.cpp b/unified-runtime/source/adapters/offload/ur_interface_loader.cpp index 17e5a74679511..da8e16f4f6a6a 100644 --- a/unified-runtime/source/adapters/offload/ur_interface_loader.cpp +++ b/unified-runtime/source/adapters/offload/ur_interface_loader.cpp @@ -388,6 +388,22 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetEnqueueExpProcAddrTable( return UR_RESULT_SUCCESS; } +UR_DLLEXPORT ur_result_t UR_APICALL urGetIPCExpProcAddrTable( + ur_api_version_t version, ur_ipc_exp_dditable_t *pDdiTable) { + auto result = validateProcInputs(version, pDdiTable); + if (UR_RESULT_SUCCESS != result) { + return result; + } + + pDdiTable->pfnGetMemHandleExp = urIPCGetMemHandleExp; + pDdiTable->pfnPutMemHandleExp = urIPCPutMemHandleExp; + pDdiTable->pfnOpenMemHandleExp = urIPCOpenMemHandleExp; + pDdiTable->pfnCloseMemHandleExp = urIPCCloseMemHandleExp; + pDdiTable->pfnGetMemHandleDataExp = urIPCGetMemHandleDataExp; + + return UR_RESULT_SUCCESS; +} + UR_DLLEXPORT ur_result_t UR_APICALL urGetProgramExpProcAddrTable( ur_api_version_t version, ur_program_exp_dditable_t *pDdiTable) { auto result = validateProcInputs(version, pDdiTable); @@ -410,6 +426,7 @@ UR_DLLEXPORT ur_result_t UR_APICALL urAllAddrTable(ur_api_version_t version, urGetContextProcAddrTable(version, &pDdiTable->Context); urGetEnqueueProcAddrTable(version, &pDdiTable->Enqueue); urGetEnqueueExpProcAddrTable(version, &pDdiTable->EnqueueExp); + urGetIPCExpProcAddrTable(version, &pDdiTable->IPCExp); urGetEventProcAddrTable(version, &pDdiTable->Event); urGetKernelProcAddrTable(version, &pDdiTable->Kernel); urGetMemProcAddrTable(version, &pDdiTable->Mem); diff --git a/unified-runtime/source/adapters/opencl/device.cpp b/unified-runtime/source/adapters/opencl/device.cpp index 4f697b05b5c88..d20f17a150aa6 100644 --- a/unified-runtime/source/adapters/opencl/device.cpp +++ b/unified-runtime/source/adapters/opencl/device.cpp @@ -1368,6 +1368,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, return ReturnValue(UR_EXP_DEVICE_2D_BLOCK_ARRAY_CAPABILITY_FLAG_LOAD | UR_EXP_DEVICE_2D_BLOCK_ARRAY_CAPABILITY_FLAG_STORE); } + case UR_DEVICE_INFO_IPC_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 19e9509987825..77d9c4a3316dd 100644 --- a/unified-runtime/source/adapters/opencl/memory.cpp +++ b/unified-runtime/source/adapters/opencl/memory.cpp @@ -579,3 +579,30 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemRelease(ur_mem_handle_t hMem) { } return UR_RESULT_SUCCESS; } + +UR_APIEXPORT ur_result_t UR_APICALL +urIPCGetMemHandleExp(ur_context_handle_t, void *, ur_exp_ipc_mem_handle_t *) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urIPCPutMemHandleExp(ur_context_handle_t, ur_exp_ipc_mem_handle_t) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +UR_APIEXPORT ur_result_t UR_APICALL urIPCOpenMemHandleExp(ur_context_handle_t, + ur_device_handle_t, + void *, size_t, + void **) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +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 urIPCGetMemHandleDataExp( + ur_context_handle_t, ur_exp_ipc_mem_handle_t, void **, size_t *) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} diff --git a/unified-runtime/source/adapters/opencl/memory.hpp b/unified-runtime/source/adapters/opencl/memory.hpp index 847ffafa76021..cc2a868a86a27 100644 --- a/unified-runtime/source/adapters/opencl/memory.hpp +++ b/unified-runtime/source/adapters/opencl/memory.hpp @@ -38,3 +38,6 @@ struct ur_mem_handle_t_ : ur::opencl::handle_base { ur_context_handle_t Ctx, ur_mem_handle_t &Mem); }; + +// IPC is currently not supported in the OpenCL adaptor. +struct ur_exp_ipc_mem_handle_t_ {}; diff --git a/unified-runtime/source/adapters/opencl/ur_interface_loader.cpp b/unified-runtime/source/adapters/opencl/ur_interface_loader.cpp index c619fa36b1ab0..18585431b7fad 100644 --- a/unified-runtime/source/adapters/opencl/ur_interface_loader.cpp +++ b/unified-runtime/source/adapters/opencl/ur_interface_loader.cpp @@ -438,6 +438,22 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetEnqueueExpProcAddrTable( return UR_RESULT_SUCCESS; } +UR_DLLEXPORT ur_result_t UR_APICALL urGetIPCExpProcAddrTable( + ur_api_version_t version, ur_ipc_exp_dditable_t *pDdiTable) { + auto result = validateProcInputs(version, pDdiTable); + if (UR_RESULT_SUCCESS != result) { + return result; + } + + pDdiTable->pfnGetMemHandleExp = urIPCGetMemHandleExp; + pDdiTable->pfnPutMemHandleExp = urIPCPutMemHandleExp; + pDdiTable->pfnOpenMemHandleExp = urIPCOpenMemHandleExp; + pDdiTable->pfnCloseMemHandleExp = urIPCCloseMemHandleExp; + pDdiTable->pfnGetMemHandleDataExp = urIPCGetMemHandleDataExp; + + return UR_RESULT_SUCCESS; +} + UR_DLLEXPORT ur_result_t UR_APICALL urGetProgramExpProcAddrTable( ur_api_version_t version, ur_program_exp_dditable_t *pDdiTable) { auto result = validateProcInputs(version, pDdiTable); @@ -460,6 +476,7 @@ UR_DLLEXPORT ur_result_t UR_APICALL urAllAddrTable(ur_api_version_t version, urGetContextProcAddrTable(version, &pDdiTable->Context); urGetEnqueueProcAddrTable(version, &pDdiTable->Enqueue); urGetEnqueueExpProcAddrTable(version, &pDdiTable->EnqueueExp); + urGetIPCExpProcAddrTable(version, &pDdiTable->IPCExp); urGetEventProcAddrTable(version, &pDdiTable->Event); urGetKernelProcAddrTable(version, &pDdiTable->Kernel); urGetMemProcAddrTable(version, &pDdiTable->Mem); diff --git a/unified-runtime/source/loader/layers/tracing/ur_trcddi.cpp b/unified-runtime/source/loader/layers/tracing/ur_trcddi.cpp index d096d3895c385..6e6929ddfca4e 100644 --- a/unified-runtime/source/loader/layers/tracing/ur_trcddi.cpp +++ b/unified-runtime/source/loader/layers/tracing/ur_trcddi.cpp @@ -9574,6 +9574,204 @@ __urdlllocal ur_result_t UR_APICALL urEnqueueTimestampRecordingExp( return result; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCGetMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCGetMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] pointer to device USM memory + void *pMem, + /// [out][alloc] pointer to the resulting IPC memory handle + ur_exp_ipc_mem_handle_t *phIPCMem) { + auto pfnGetMemHandleExp = getContext()->urDdiTable.IPCExp.pfnGetMemHandleExp; + + if (nullptr == pfnGetMemHandleExp) + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + + ur_ipc_get_mem_handle_exp_params_t params = {&hContext, &pMem, &phIPCMem}; + uint64_t instance = getContext()->notify_begin( + UR_FUNCTION_IPC_GET_MEM_HANDLE_EXP, "urIPCGetMemHandleExp", ¶ms); + + auto &logger = getContext()->logger; + UR_LOG_L(logger, INFO, " ---> urIPCGetMemHandleExp\n"); + + ur_result_t result = pfnGetMemHandleExp(hContext, pMem, phIPCMem); + + getContext()->notify_end(UR_FUNCTION_IPC_GET_MEM_HANDLE_EXP, + "urIPCGetMemHandleExp", ¶ms, &result, instance); + + if (logger.getLevel() <= UR_LOGGER_LEVEL_INFO) { + std::ostringstream args_str; + ur::extras::printFunctionParams( + args_str, UR_FUNCTION_IPC_GET_MEM_HANDLE_EXP, ¶ms); + UR_LOG_L(logger, INFO, " <--- urIPCGetMemHandleExp({}) -> {};\n", + args_str.str(), result); + } + + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCPutMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCPutMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] the IPC memory handle + ur_exp_ipc_mem_handle_t hIPCMem) { + auto pfnPutMemHandleExp = getContext()->urDdiTable.IPCExp.pfnPutMemHandleExp; + + if (nullptr == pfnPutMemHandleExp) + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + + ur_ipc_put_mem_handle_exp_params_t params = {&hContext, &hIPCMem}; + uint64_t instance = getContext()->notify_begin( + UR_FUNCTION_IPC_PUT_MEM_HANDLE_EXP, "urIPCPutMemHandleExp", ¶ms); + + auto &logger = getContext()->logger; + UR_LOG_L(logger, INFO, " ---> urIPCPutMemHandleExp\n"); + + ur_result_t result = pfnPutMemHandleExp(hContext, hIPCMem); + + getContext()->notify_end(UR_FUNCTION_IPC_PUT_MEM_HANDLE_EXP, + "urIPCPutMemHandleExp", ¶ms, &result, instance); + + if (logger.getLevel() <= UR_LOGGER_LEVEL_INFO) { + std::ostringstream args_str; + ur::extras::printFunctionParams( + args_str, UR_FUNCTION_IPC_PUT_MEM_HANDLE_EXP, ¶ms); + UR_LOG_L(logger, INFO, " <--- urIPCPutMemHandleExp({}) -> {};\n", + args_str.str(), result); + } + + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCOpenMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCOpenMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object the corresponding USM device memory + /// was allocated on + ur_device_handle_t hDevice, + /// [in] the IPC memory handle data + void *ipcMemHandleData, + /// [in] size of the IPC memory handle data + size_t ipcMemHandleDataSize, + /// [out] pointer to a pointer to device USM memory + void **ppMem) { + auto pfnOpenMemHandleExp = + getContext()->urDdiTable.IPCExp.pfnOpenMemHandleExp; + + if (nullptr == pfnOpenMemHandleExp) + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + + ur_ipc_open_mem_handle_exp_params_t params = { + &hContext, &hDevice, &ipcMemHandleData, &ipcMemHandleDataSize, &ppMem}; + uint64_t instance = getContext()->notify_begin( + UR_FUNCTION_IPC_OPEN_MEM_HANDLE_EXP, "urIPCOpenMemHandleExp", ¶ms); + + auto &logger = getContext()->logger; + UR_LOG_L(logger, INFO, " ---> urIPCOpenMemHandleExp\n"); + + ur_result_t result = pfnOpenMemHandleExp(hContext, hDevice, ipcMemHandleData, + ipcMemHandleDataSize, ppMem); + + getContext()->notify_end(UR_FUNCTION_IPC_OPEN_MEM_HANDLE_EXP, + "urIPCOpenMemHandleExp", ¶ms, &result, instance); + + if (logger.getLevel() <= UR_LOGGER_LEVEL_INFO) { + std::ostringstream args_str; + ur::extras::printFunctionParams( + args_str, UR_FUNCTION_IPC_OPEN_MEM_HANDLE_EXP, ¶ms); + UR_LOG_L(logger, INFO, " <--- urIPCOpenMemHandleExp({}) -> {};\n", + args_str.str(), result); + } + + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCCloseMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCCloseMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] pointer to device USM memory opened through urIPCOpenMemHandleExp + void *pMem) { + auto pfnCloseMemHandleExp = + getContext()->urDdiTable.IPCExp.pfnCloseMemHandleExp; + + if (nullptr == pfnCloseMemHandleExp) + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + + ur_ipc_close_mem_handle_exp_params_t params = {&hContext, &pMem}; + uint64_t instance = getContext()->notify_begin( + UR_FUNCTION_IPC_CLOSE_MEM_HANDLE_EXP, "urIPCCloseMemHandleExp", ¶ms); + + auto &logger = getContext()->logger; + UR_LOG_L(logger, INFO, " ---> urIPCCloseMemHandleExp\n"); + + ur_result_t result = pfnCloseMemHandleExp(hContext, pMem); + + getContext()->notify_end(UR_FUNCTION_IPC_CLOSE_MEM_HANDLE_EXP, + "urIPCCloseMemHandleExp", ¶ms, &result, + instance); + + if (logger.getLevel() <= UR_LOGGER_LEVEL_INFO) { + std::ostringstream args_str; + ur::extras::printFunctionParams( + args_str, UR_FUNCTION_IPC_CLOSE_MEM_HANDLE_EXP, ¶ms); + UR_LOG_L(logger, INFO, " <--- urIPCCloseMemHandleExp({}) -> {};\n", + args_str.str(), result); + } + + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCGetMemHandleDataExp +__urdlllocal ur_result_t UR_APICALL urIPCGetMemHandleDataExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] the IPC memory handle + ur_exp_ipc_mem_handle_t hIPCMem, + /// [out][optional] a pointer to the IPC memory handle data + void **ppIPCHandleData, + /// [out][optional] size of the resulting IPC memory handle data + size_t *pIPCMemHandleDataSizeRet) { + auto pfnGetMemHandleDataExp = + getContext()->urDdiTable.IPCExp.pfnGetMemHandleDataExp; + + if (nullptr == pfnGetMemHandleDataExp) + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + + ur_ipc_get_mem_handle_data_exp_params_t params = { + &hContext, &hIPCMem, &ppIPCHandleData, &pIPCMemHandleDataSizeRet}; + uint64_t instance = + getContext()->notify_begin(UR_FUNCTION_IPC_GET_MEM_HANDLE_DATA_EXP, + "urIPCGetMemHandleDataExp", ¶ms); + + auto &logger = getContext()->logger; + UR_LOG_L(logger, INFO, " ---> urIPCGetMemHandleDataExp\n"); + + ur_result_t result = pfnGetMemHandleDataExp( + hContext, hIPCMem, ppIPCHandleData, pIPCMemHandleDataSizeRet); + + getContext()->notify_end(UR_FUNCTION_IPC_GET_MEM_HANDLE_DATA_EXP, + "urIPCGetMemHandleDataExp", ¶ms, &result, + instance); + + if (logger.getLevel() <= UR_LOGGER_LEVEL_INFO) { + std::ostringstream args_str; + ur::extras::printFunctionParams( + args_str, UR_FUNCTION_IPC_GET_MEM_HANDLE_DATA_EXP, ¶ms); + UR_LOG_L(logger, INFO, " <--- urIPCGetMemHandleDataExp({}) -> {};\n", + args_str.str(), result); + } + + return result; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Intercept function for urMemoryExportAllocExportableMemoryExp __urdlllocal ur_result_t UR_APICALL urMemoryExportAllocExportableMemoryExp( @@ -10783,6 +10981,50 @@ __urdlllocal ur_result_t UR_APICALL urGetEventProcAddrTable( return result; } /////////////////////////////////////////////////////////////////////////////// +/// @brief Exported function for filling application's IPCExp table +/// with current process' addresses +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// - ::UR_RESULT_ERROR_UNSUPPORTED_VERSION +__urdlllocal ur_result_t UR_APICALL urGetIPCExpProcAddrTable( + /// [in] API version requested + ur_api_version_t version, + /// [in,out] pointer to table of DDI function pointers + ur_ipc_exp_dditable_t *pDdiTable) { + auto &dditable = ur_tracing_layer::getContext()->urDdiTable.IPCExp; + + if (nullptr == pDdiTable) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + + if (UR_MAJOR_VERSION(ur_tracing_layer::getContext()->version) != + UR_MAJOR_VERSION(version) || + UR_MINOR_VERSION(ur_tracing_layer::getContext()->version) > + UR_MINOR_VERSION(version)) + return UR_RESULT_ERROR_UNSUPPORTED_VERSION; + + ur_result_t result = UR_RESULT_SUCCESS; + + dditable.pfnGetMemHandleExp = pDdiTable->pfnGetMemHandleExp; + pDdiTable->pfnGetMemHandleExp = ur_tracing_layer::urIPCGetMemHandleExp; + + dditable.pfnPutMemHandleExp = pDdiTable->pfnPutMemHandleExp; + pDdiTable->pfnPutMemHandleExp = ur_tracing_layer::urIPCPutMemHandleExp; + + dditable.pfnOpenMemHandleExp = pDdiTable->pfnOpenMemHandleExp; + pDdiTable->pfnOpenMemHandleExp = ur_tracing_layer::urIPCOpenMemHandleExp; + + dditable.pfnCloseMemHandleExp = pDdiTable->pfnCloseMemHandleExp; + pDdiTable->pfnCloseMemHandleExp = ur_tracing_layer::urIPCCloseMemHandleExp; + + dditable.pfnGetMemHandleDataExp = pDdiTable->pfnGetMemHandleDataExp; + pDdiTable->pfnGetMemHandleDataExp = + ur_tracing_layer::urIPCGetMemHandleDataExp; + + return result; +} +/////////////////////////////////////////////////////////////////////////////// /// @brief Exported function for filling application's Kernel table /// with current process' addresses /// @@ -11590,6 +11832,11 @@ ur_result_t context_t::init(ur_dditable_t *dditable, &dditable->Event); } + if (UR_RESULT_SUCCESS == result) { + result = ur_tracing_layer::urGetIPCExpProcAddrTable(UR_API_VERSION_CURRENT, + &dditable->IPCExp); + } + if (UR_RESULT_SUCCESS == result) { result = ur_tracing_layer::urGetKernelProcAddrTable(UR_API_VERSION_CURRENT, &dditable->Kernel); diff --git a/unified-runtime/source/loader/layers/validation/ur_valddi.cpp b/unified-runtime/source/loader/layers/validation/ur_valddi.cpp index 32dec6f1b25df..e176955907a92 100644 --- a/unified-runtime/source/loader/layers/validation/ur_valddi.cpp +++ b/unified-runtime/source/loader/layers/validation/ur_valddi.cpp @@ -10337,6 +10337,190 @@ __urdlllocal ur_result_t UR_APICALL urEnqueueTimestampRecordingExp( return result; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCGetMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCGetMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] pointer to device USM memory + void *pMem, + /// [out][alloc] pointer to the resulting IPC memory handle + ur_exp_ipc_mem_handle_t *phIPCMem) { + auto pfnGetMemHandleExp = getContext()->urDdiTable.IPCExp.pfnGetMemHandleExp; + + if (nullptr == pfnGetMemHandleExp) { + return UR_RESULT_ERROR_UNINITIALIZED; + } + + if (getContext()->enableParameterValidation) { + if (NULL == phIPCMem) + 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 = pfnGetMemHandleExp(hContext, pMem, phIPCMem); + + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCPutMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCPutMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] the IPC memory handle + ur_exp_ipc_mem_handle_t hIPCMem) { + auto pfnPutMemHandleExp = getContext()->urDdiTable.IPCExp.pfnPutMemHandleExp; + + if (nullptr == pfnPutMemHandleExp) { + return UR_RESULT_ERROR_UNINITIALIZED; + } + + if (getContext()->enableParameterValidation) { + if (NULL == hContext) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + + if (NULL == hIPCMem) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + } + + if (getContext()->enableLifetimeValidation && + !getContext()->refCountContext->isReferenceValid(hContext)) { + URLOG_CTX_INVALID_REFERENCE(hContext); + } + + ur_result_t result = pfnPutMemHandleExp(hContext, hIPCMem); + + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCOpenMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCOpenMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object the corresponding USM device memory + /// was allocated on + ur_device_handle_t hDevice, + /// [in] the IPC memory handle data + void *ipcMemHandleData, + /// [in] size of the IPC memory handle data + size_t ipcMemHandleDataSize, + /// [out] pointer to a pointer to device USM memory + void **ppMem) { + auto pfnOpenMemHandleExp = + getContext()->urDdiTable.IPCExp.pfnOpenMemHandleExp; + + if (nullptr == pfnOpenMemHandleExp) { + return UR_RESULT_ERROR_UNINITIALIZED; + } + + if (getContext()->enableParameterValidation) { + if (NULL == ipcMemHandleData) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + + if (NULL == ppMem) + 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 = pfnOpenMemHandleExp(hContext, hDevice, ipcMemHandleData, + ipcMemHandleDataSize, ppMem); + + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCCloseMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCCloseMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] pointer to device USM memory opened through urIPCOpenMemHandleExp + void *pMem) { + auto pfnCloseMemHandleExp = + getContext()->urDdiTable.IPCExp.pfnCloseMemHandleExp; + + if (nullptr == pfnCloseMemHandleExp) { + return UR_RESULT_ERROR_UNINITIALIZED; + } + + if (getContext()->enableParameterValidation) { + if (NULL == pMem) + 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 = pfnCloseMemHandleExp(hContext, pMem); + + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCGetMemHandleDataExp +__urdlllocal ur_result_t UR_APICALL urIPCGetMemHandleDataExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] the IPC memory handle + ur_exp_ipc_mem_handle_t hIPCMem, + /// [out][optional] a pointer to the IPC memory handle data + void **ppIPCHandleData, + /// [out][optional] size of the resulting IPC memory handle data + size_t *pIPCMemHandleDataSizeRet) { + auto pfnGetMemHandleDataExp = + getContext()->urDdiTable.IPCExp.pfnGetMemHandleDataExp; + + if (nullptr == pfnGetMemHandleDataExp) { + return UR_RESULT_ERROR_UNINITIALIZED; + } + + if (getContext()->enableParameterValidation) { + if (NULL == hContext) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + + if (NULL == hIPCMem) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + } + + if (getContext()->enableLifetimeValidation && + !getContext()->refCountContext->isReferenceValid(hContext)) { + URLOG_CTX_INVALID_REFERENCE(hContext); + } + + ur_result_t result = pfnGetMemHandleDataExp( + hContext, hIPCMem, ppIPCHandleData, pIPCMemHandleDataSizeRet); + + return result; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Intercept function for urMemoryExportAllocExportableMemoryExp __urdlllocal ur_result_t UR_APICALL urMemoryExportAllocExportableMemoryExp( @@ -11585,6 +11769,51 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetEventProcAddrTable( return result; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Exported function for filling application's IPCExp table +/// with current process' addresses +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// - ::UR_RESULT_ERROR_UNSUPPORTED_VERSION +UR_DLLEXPORT ur_result_t UR_APICALL urGetIPCExpProcAddrTable( + /// [in] API version requested + ur_api_version_t version, + /// [in,out] pointer to table of DDI function pointers + ur_ipc_exp_dditable_t *pDdiTable) { + auto &dditable = ur_validation_layer::getContext()->urDdiTable.IPCExp; + + if (nullptr == pDdiTable) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + + if (UR_MAJOR_VERSION(ur_validation_layer::getContext()->version) != + UR_MAJOR_VERSION(version) || + UR_MINOR_VERSION(ur_validation_layer::getContext()->version) > + UR_MINOR_VERSION(version)) + return UR_RESULT_ERROR_UNSUPPORTED_VERSION; + + ur_result_t result = UR_RESULT_SUCCESS; + + dditable.pfnGetMemHandleExp = pDdiTable->pfnGetMemHandleExp; + pDdiTable->pfnGetMemHandleExp = ur_validation_layer::urIPCGetMemHandleExp; + + dditable.pfnPutMemHandleExp = pDdiTable->pfnPutMemHandleExp; + pDdiTable->pfnPutMemHandleExp = ur_validation_layer::urIPCPutMemHandleExp; + + dditable.pfnOpenMemHandleExp = pDdiTable->pfnOpenMemHandleExp; + pDdiTable->pfnOpenMemHandleExp = ur_validation_layer::urIPCOpenMemHandleExp; + + dditable.pfnCloseMemHandleExp = pDdiTable->pfnCloseMemHandleExp; + pDdiTable->pfnCloseMemHandleExp = ur_validation_layer::urIPCCloseMemHandleExp; + + dditable.pfnGetMemHandleDataExp = pDdiTable->pfnGetMemHandleDataExp; + pDdiTable->pfnGetMemHandleDataExp = + ur_validation_layer::urIPCGetMemHandleDataExp; + + return result; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Exported function for filling application's Kernel table /// with current process' addresses @@ -12427,6 +12656,11 @@ ur_result_t context_t::init(ur_dditable_t *dditable, UR_API_VERSION_CURRENT, &dditable->Event); } + if (UR_RESULT_SUCCESS == result) { + result = ur_validation_layer::urGetIPCExpProcAddrTable( + UR_API_VERSION_CURRENT, &dditable->IPCExp); + } + if (UR_RESULT_SUCCESS == result) { result = ur_validation_layer::urGetKernelProcAddrTable( UR_API_VERSION_CURRENT, &dditable->Kernel); diff --git a/unified-runtime/source/loader/loader.def.in b/unified-runtime/source/loader/loader.def.in index 3ad47149315ee..a09ecdf084ddd 100644 --- a/unified-runtime/source/loader/loader.def.in +++ b/unified-runtime/source/loader/loader.def.in @@ -118,6 +118,7 @@ EXPORTS urGetEnqueueExpProcAddrTable urGetEnqueueProcAddrTable urGetEventProcAddrTable + urGetIPCExpProcAddrTable urGetKernelProcAddrTable urGetMemProcAddrTable urGetMemoryExportExpProcAddrTable @@ -131,6 +132,11 @@ EXPORTS urGetUSMProcAddrTable urGetUsmP2PExpProcAddrTable urGetVirtualMemProcAddrTable + urIPCCloseMemHandleExp + urIPCGetMemHandleDataExp + urIPCGetMemHandleExp + urIPCOpenMemHandleExp + urIPCPutMemHandleExp urKernelCreate urKernelCreateWithNativeHandle urKernelGetGroupInfo @@ -361,6 +367,11 @@ EXPORTS urPrintImageDesc urPrintImageFormat urPrintImageInfo + urPrintIpcCloseMemHandleExpParams + urPrintIpcGetMemHandleDataExpParams + urPrintIpcGetMemHandleExpParams + urPrintIpcOpenMemHandleExpParams + urPrintIpcPutMemHandleExpParams urPrintKernelArgLocalProperties urPrintKernelArgMemObjProperties urPrintKernelArgPointerProperties diff --git a/unified-runtime/source/loader/loader.map.in b/unified-runtime/source/loader/loader.map.in index fde803f9aa45a..08c16c8e43c42 100644 --- a/unified-runtime/source/loader/loader.map.in +++ b/unified-runtime/source/loader/loader.map.in @@ -118,6 +118,7 @@ urGetEnqueueExpProcAddrTable; urGetEnqueueProcAddrTable; urGetEventProcAddrTable; + urGetIPCExpProcAddrTable; urGetKernelProcAddrTable; urGetMemProcAddrTable; urGetMemoryExportExpProcAddrTable; @@ -131,6 +132,11 @@ urGetUSMProcAddrTable; urGetUsmP2PExpProcAddrTable; urGetVirtualMemProcAddrTable; + urIPCCloseMemHandleExp; + urIPCGetMemHandleDataExp; + urIPCGetMemHandleExp; + urIPCOpenMemHandleExp; + urIPCPutMemHandleExp; urKernelCreate; urKernelCreateWithNativeHandle; urKernelGetGroupInfo; @@ -361,6 +367,11 @@ urPrintImageDesc; urPrintImageFormat; urPrintImageInfo; + urPrintIpcCloseMemHandleExpParams; + urPrintIpcGetMemHandleDataExpParams; + urPrintIpcGetMemHandleExpParams; + urPrintIpcOpenMemHandleExpParams; + urPrintIpcPutMemHandleExpParams; urPrintKernelArgLocalProperties; urPrintKernelArgMemObjProperties; urPrintKernelArgPointerProperties; diff --git a/unified-runtime/source/loader/ur_ldrddi.cpp b/unified-runtime/source/loader/ur_ldrddi.cpp index 75ae04bc5a4a8..26b7b74e0f528 100644 --- a/unified-runtime/source/loader/ur_ldrddi.cpp +++ b/unified-runtime/source/loader/ur_ldrddi.cpp @@ -5452,6 +5452,111 @@ __urdlllocal ur_result_t UR_APICALL urEnqueueTimestampRecordingExp( phEventWaitList, phEvent); } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCGetMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCGetMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] pointer to device USM memory + void *pMem, + /// [out][alloc] pointer to the resulting IPC memory handle + ur_exp_ipc_mem_handle_t *phIPCMem) { + + auto *dditable = *reinterpret_cast(hContext); + + auto *pfnGetMemHandleExp = dditable->IPCExp.pfnGetMemHandleExp; + if (nullptr == pfnGetMemHandleExp) + return UR_RESULT_ERROR_UNINITIALIZED; + + // forward to device-platform + return pfnGetMemHandleExp(hContext, pMem, phIPCMem); +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCPutMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCPutMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] the IPC memory handle + ur_exp_ipc_mem_handle_t hIPCMem) { + + auto *dditable = *reinterpret_cast(hContext); + + auto *pfnPutMemHandleExp = dditable->IPCExp.pfnPutMemHandleExp; + if (nullptr == pfnPutMemHandleExp) + return UR_RESULT_ERROR_UNINITIALIZED; + + // forward to device-platform + return pfnPutMemHandleExp(hContext, hIPCMem); +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCOpenMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCOpenMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object the corresponding USM device memory + /// was allocated on + ur_device_handle_t hDevice, + /// [in] the IPC memory handle data + void *ipcMemHandleData, + /// [in] size of the IPC memory handle data + size_t ipcMemHandleDataSize, + /// [out] pointer to a pointer to device USM memory + void **ppMem) { + + auto *dditable = *reinterpret_cast(hContext); + + auto *pfnOpenMemHandleExp = dditable->IPCExp.pfnOpenMemHandleExp; + if (nullptr == pfnOpenMemHandleExp) + return UR_RESULT_ERROR_UNINITIALIZED; + + // forward to device-platform + return pfnOpenMemHandleExp(hContext, hDevice, ipcMemHandleData, + ipcMemHandleDataSize, ppMem); +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCCloseMemHandleExp +__urdlllocal ur_result_t UR_APICALL urIPCCloseMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] pointer to device USM memory opened through urIPCOpenMemHandleExp + void *pMem) { + + auto *dditable = *reinterpret_cast(hContext); + + auto *pfnCloseMemHandleExp = dditable->IPCExp.pfnCloseMemHandleExp; + if (nullptr == pfnCloseMemHandleExp) + return UR_RESULT_ERROR_UNINITIALIZED; + + // forward to device-platform + return pfnCloseMemHandleExp(hContext, pMem); +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for urIPCGetMemHandleDataExp +__urdlllocal ur_result_t UR_APICALL urIPCGetMemHandleDataExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] the IPC memory handle + ur_exp_ipc_mem_handle_t hIPCMem, + /// [out][optional] a pointer to the IPC memory handle data + void **ppIPCHandleData, + /// [out][optional] size of the resulting IPC memory handle data + size_t *pIPCMemHandleDataSizeRet) { + + auto *dditable = *reinterpret_cast(hContext); + + auto *pfnGetMemHandleDataExp = dditable->IPCExp.pfnGetMemHandleDataExp; + if (nullptr == pfnGetMemHandleDataExp) + return UR_RESULT_ERROR_UNINITIALIZED; + + // forward to device-platform + return pfnGetMemHandleDataExp(hContext, hIPCMem, ppIPCHandleData, + pIPCMemHandleDataSizeRet); +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Intercept function for urMemoryExportAllocExportableMemoryExp __urdlllocal ur_result_t UR_APICALL urMemoryExportAllocExportableMemoryExp( @@ -6322,6 +6427,62 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetEventProcAddrTable( return result; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Exported function for filling application's IPCExp table +/// with current process' addresses +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// - ::UR_RESULT_ERROR_UNSUPPORTED_VERSION +UR_DLLEXPORT ur_result_t UR_APICALL urGetIPCExpProcAddrTable( + /// [in] API version requested + ur_api_version_t version, + /// [in,out] pointer to table of DDI function pointers + ur_ipc_exp_dditable_t *pDdiTable) { + if (nullptr == pDdiTable) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + + if (ur_loader::getContext()->version < version) + return UR_RESULT_ERROR_UNSUPPORTED_VERSION; + + ur_result_t result = UR_RESULT_SUCCESS; + + // Load the device-platform DDI tables + for (auto &platform : ur_loader::getContext()->platforms) { + // statically linked adapter inside of the loader + if (platform.handle == nullptr) + continue; + + if (platform.initStatus != UR_RESULT_SUCCESS) + continue; + auto getTable = reinterpret_cast( + ur_loader::LibLoader::getFunctionPtr(platform.handle.get(), + "urGetIPCExpProcAddrTable")); + if (!getTable) + continue; + platform.initStatus = getTable(version, &platform.dditable.IPCExp); + } + + if (UR_RESULT_SUCCESS == result) { + if (ur_loader::getContext()->platforms.size() != 1 || + ur_loader::getContext()->forceIntercept) { + // return pointers to loader's DDIs + pDdiTable->pfnGetMemHandleExp = ur_loader::urIPCGetMemHandleExp; + pDdiTable->pfnPutMemHandleExp = ur_loader::urIPCPutMemHandleExp; + pDdiTable->pfnOpenMemHandleExp = ur_loader::urIPCOpenMemHandleExp; + pDdiTable->pfnCloseMemHandleExp = ur_loader::urIPCCloseMemHandleExp; + pDdiTable->pfnGetMemHandleDataExp = ur_loader::urIPCGetMemHandleDataExp; + } else { + // return pointers directly to platform's DDIs + *pDdiTable = ur_loader::getContext()->platforms.front().dditable.IPCExp; + } + } + + return result; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Exported function for filling application's Kernel table /// with current process' addresses diff --git a/unified-runtime/source/loader/ur_libapi.cpp b/unified-runtime/source/loader/ur_libapi.cpp index 4ec2282647e80..0d49101533af8 100644 --- a/unified-runtime/source/loader/ur_libapi.cpp +++ b/unified-runtime/source/loader/ur_libapi.cpp @@ -9997,6 +9997,176 @@ ur_result_t UR_APICALL urEnqueueTimestampRecordingExp( return exceptionToResult(std::current_exception()); } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Gets an inter-process memory handle for a pointer to device USM +/// memory +/// +/// @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 == phIPCMem` +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +ur_result_t UR_APICALL urIPCGetMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] pointer to device USM memory + void *pMem, + /// [out][alloc] pointer to the resulting IPC memory handle + ur_exp_ipc_mem_handle_t *phIPCMem) try { + auto pfnGetMemHandleExp = + ur_lib::getContext()->urDdiTable.IPCExp.pfnGetMemHandleExp; + if (nullptr == pfnGetMemHandleExp) + return UR_RESULT_ERROR_UNINITIALIZED; + + return pfnGetMemHandleExp(hContext, pMem, phIPCMem); +} catch (...) { + return exceptionToResult(std::current_exception()); +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Releases an inter-process 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 == hIPCMem` +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +ur_result_t UR_APICALL urIPCPutMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] the IPC memory handle + ur_exp_ipc_mem_handle_t hIPCMem) try { + auto pfnPutMemHandleExp = + ur_lib::getContext()->urDdiTable.IPCExp.pfnPutMemHandleExp; + if (nullptr == pfnPutMemHandleExp) + return UR_RESULT_ERROR_UNINITIALIZED; + + return pfnPutMemHandleExp(hContext, hIPCMem); +} catch (...) { + return exceptionToResult(std::current_exception()); +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Opens an inter-process memory handle from raw data to get the +/// corresponding pointer to device USM memory +/// +/// @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_CONTEXT +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == ipcMemHandleData` +/// + `NULL == ppMem` +/// - ::UR_RESULT_ERROR_INVALID_VALUE +/// + ipcMemHandleDataSize is not the same as the size of IPC memory +/// handle data +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +ur_result_t UR_APICALL urIPCOpenMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object the corresponding USM device memory + /// was allocated on + ur_device_handle_t hDevice, + /// [in] the IPC memory handle data + void *ipcMemHandleData, + /// [in] size of the IPC memory handle data + size_t ipcMemHandleDataSize, + /// [out] pointer to a pointer to device USM memory + void **ppMem) try { + auto pfnOpenMemHandleExp = + ur_lib::getContext()->urDdiTable.IPCExp.pfnOpenMemHandleExp; + if (nullptr == pfnOpenMemHandleExp) + return UR_RESULT_ERROR_UNINITIALIZED; + + return pfnOpenMemHandleExp(hContext, hDevice, ipcMemHandleData, + ipcMemHandleDataSize, ppMem); +} catch (...) { + return exceptionToResult(std::current_exception()); +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Closes an inter-process 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_CONTEXT +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pMem` +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +ur_result_t UR_APICALL urIPCCloseMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] pointer to device USM memory opened through urIPCOpenMemHandleExp + void *pMem) try { + auto pfnCloseMemHandleExp = + ur_lib::getContext()->urDdiTable.IPCExp.pfnCloseMemHandleExp; + if (nullptr == pfnCloseMemHandleExp) + return UR_RESULT_ERROR_UNINITIALIZED; + + return pfnCloseMemHandleExp(hContext, pMem); +} catch (...) { + return exceptionToResult(std::current_exception()); +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Gets the data of an inter-process 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 == hIPCMem` +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +ur_result_t UR_APICALL urIPCGetMemHandleDataExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] the IPC memory handle + ur_exp_ipc_mem_handle_t hIPCMem, + /// [out][optional] a pointer to the IPC memory handle data + void **ppIPCHandleData, + /// [out][optional] size of the resulting IPC memory handle data + size_t *pIPCMemHandleDataSizeRet) try { + auto pfnGetMemHandleDataExp = + ur_lib::getContext()->urDdiTable.IPCExp.pfnGetMemHandleDataExp; + if (nullptr == pfnGetMemHandleDataExp) + return UR_RESULT_ERROR_UNINITIALIZED; + + return pfnGetMemHandleDataExp(hContext, hIPCMem, ppIPCHandleData, + pIPCMemHandleDataSizeRet); +} 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_libddi.cpp b/unified-runtime/source/loader/ur_libddi.cpp index b06ee531cdb46..648fc87872188 100644 --- a/unified-runtime/source/loader/ur_libddi.cpp +++ b/unified-runtime/source/loader/ur_libddi.cpp @@ -56,6 +56,11 @@ __urdlllocal ur_result_t context_t::ddiInit() { result = urGetEventProcAddrTable(UR_API_VERSION_CURRENT, &urDdiTable.Event); } + if (UR_RESULT_SUCCESS == result) { + result = + urGetIPCExpProcAddrTable(UR_API_VERSION_CURRENT, &urDdiTable.IPCExp); + } + if (UR_RESULT_SUCCESS == result) { result = urGetKernelProcAddrTable(UR_API_VERSION_CURRENT, &urDdiTable.Kernel); diff --git a/unified-runtime/source/loader/ur_print.cpp b/unified-runtime/source/loader/ur_print.cpp index f3d5c96e376ca..4c5552733bd38 100644 --- a/unified-runtime/source/loader/ur_print.cpp +++ b/unified-runtime/source/loader/ur_print.cpp @@ -2000,6 +2000,46 @@ ur_result_t urPrintEventSetCallbackParams( return str_copy(&ss, buffer, buff_size, out_size); } +ur_result_t urPrintIpcGetMemHandleExpParams( + const struct ur_ipc_get_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 urPrintIpcPutMemHandleExpParams( + const struct ur_ipc_put_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 urPrintIpcOpenMemHandleExpParams( + const struct ur_ipc_open_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 urPrintIpcCloseMemHandleExpParams( + const struct ur_ipc_close_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 urPrintIpcGetMemHandleDataExpParams( + const struct ur_ipc_get_mem_handle_data_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 8e3424b693e62..250e1af005851 100644 --- a/unified-runtime/source/ur_api.cpp +++ b/unified-runtime/source/ur_api.cpp @@ -8704,6 +8704,144 @@ ur_result_t UR_APICALL urEnqueueTimestampRecordingExp( return result; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Gets an inter-process memory handle for a pointer to device USM +/// memory +/// +/// @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 == phIPCMem` +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +ur_result_t UR_APICALL urIPCGetMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] pointer to device USM memory + void *pMem, + /// [out][alloc] pointer to the resulting IPC memory handle + ur_exp_ipc_mem_handle_t *phIPCMem) { + ur_result_t result = UR_RESULT_SUCCESS; + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Releases an inter-process 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 == hIPCMem` +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +ur_result_t UR_APICALL urIPCPutMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] the IPC memory handle + ur_exp_ipc_mem_handle_t hIPCMem) { + ur_result_t result = UR_RESULT_SUCCESS; + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Opens an inter-process memory handle from raw data to get the +/// corresponding pointer to device USM memory +/// +/// @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_CONTEXT +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == ipcMemHandleData` +/// + `NULL == ppMem` +/// - ::UR_RESULT_ERROR_INVALID_VALUE +/// + ipcMemHandleDataSize is not the same as the size of IPC memory +/// handle data +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +ur_result_t UR_APICALL urIPCOpenMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object the corresponding USM device memory + /// was allocated on + ur_device_handle_t hDevice, + /// [in] the IPC memory handle data + void *ipcMemHandleData, + /// [in] size of the IPC memory handle data + size_t ipcMemHandleDataSize, + /// [out] pointer to a pointer to device USM memory + void **ppMem) { + ur_result_t result = UR_RESULT_SUCCESS; + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Closes an inter-process 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_CONTEXT +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pMem` +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +ur_result_t UR_APICALL urIPCCloseMemHandleExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] pointer to device USM memory opened through urIPCOpenMemHandleExp + void *pMem) { + ur_result_t result = UR_RESULT_SUCCESS; + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Gets the data of an inter-process 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 == hIPCMem` +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +/// - ::UR_RESULT_ERROR_OUT_OF_HOST_MEMORY +/// - ::UR_RESULT_ERROR_OUT_OF_RESOURCES +ur_result_t UR_APICALL urIPCGetMemHandleDataExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] the IPC memory handle + ur_exp_ipc_mem_handle_t hIPCMem, + /// [out][optional] a pointer to the IPC memory handle data + void **ppIPCHandleData, + /// [out][optional] size of the resulting IPC memory handle data + size_t *pIPCMemHandleDataSizeRet) { + 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/tools/urinfo/urinfo.hpp b/unified-runtime/tools/urinfo/urinfo.hpp index b08661787cccf..d27bf7a229ff9 100644 --- a/unified-runtime/tools/urinfo/urinfo.hpp +++ b/unified-runtime/tools/urinfo/urinfo.hpp @@ -445,6 +445,8 @@ inline void printDeviceInfos(ur_device_handle_t hDevice, printDeviceInfo( hDevice, UR_DEVICE_INFO_2D_BLOCK_ARRAY_CAPABILITIES_EXP); std::cout << prefix; + printDeviceInfo(hDevice, UR_DEVICE_INFO_IPC_MEMORY_SUPPORT_EXP); + std::cout << prefix; printDeviceInfo(hDevice, UR_DEVICE_INFO_ASYNC_USM_ALLOCATIONS_SUPPORT_EXP); std::cout << prefix;