diff --git a/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td b/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td index 2c5903a033301..8e7807f9ddae0 100644 --- a/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td +++ b/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td @@ -106,6 +106,7 @@ def Aspectext_intel_eus_per_xe_core : Aspect<"ext_intel_eus_per_xe_core">; def Aspectext_intel_max_lanes_per_hw_thread : Aspect<"ext_intel_max_lanes_per_hw_thread">; def Aspectext_oneapi_ipc_physical_memory : Aspect<"ext_oneapi_ipc_physical_memory">; def Aspectext_oneapi_register_host_memory : Aspect<"ext_oneapi_register_host_memory">; +def Aspectext_oneapi_ipc_event : Aspect<"ext_oneapi_ipc_event">; // Deprecated aspects def AspectInt64_base_atomics : Aspect<"int64_base_atomics">; @@ -193,7 +194,8 @@ def : TargetInfo<"__TestAspectList", Aspectext_intel_eus_per_xe_core, Aspectext_intel_max_lanes_per_hw_thread, Aspectext_oneapi_ipc_physical_memory, - Aspectext_oneapi_register_host_memory], + Aspectext_oneapi_register_host_memory, + Aspectext_oneapi_ipc_event], []>; // 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/include/sycl/event.hpp b/sycl/include/sycl/event.hpp index d815f9819b9bf..806d16d4cb74f 100644 --- a/sycl/include/sycl/event.hpp +++ b/sycl/include/sycl/event.hpp @@ -147,6 +147,10 @@ class __SYCL_EXPORT event : public detail::OwnerLessBase { /// \return the backend associated with this platform backend get_backend() const noexcept; + /// Returns true if the event was created with inter-process sharing + /// enabled. + bool ext_oneapi_ipc_enabled() const noexcept; + private: event(std::shared_ptr EventImpl); diff --git a/sycl/include/sycl/ext/oneapi/experimental/detail/ipc_common.hpp b/sycl/include/sycl/ext/oneapi/experimental/detail/ipc_common.hpp index a16d1919b6b83..d95c9fd421ecf 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/detail/ipc_common.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/detail/ipc_common.hpp @@ -8,11 +8,33 @@ #pragma once +#include +#include + +namespace sycl { +inline namespace _V1 { + +namespace ext::oneapi::experimental { + +// Shared by the inter-process-communication extensions (physical_mem and +// event); each target adds an is_property_key_of specialisation in its own +// header. +struct enable_ipc_key + : detail::compile_time_property_key { + using value_t = property_value; +}; + +inline constexpr enable_ipc_key::value_t enable_ipc; + +} // namespace ext::oneapi::experimental +} // namespace _V1 +} // namespace sycl + #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0) #include -#include #include +#include #include #include @@ -42,6 +64,11 @@ __SYCL_EXPORT handle get(physical_mem &physmem); __SYCL_EXPORT void put(handle &ipc_handle, const sycl::context &ctx); } // namespace ext::oneapi::experimental::ipc::physical_memory +namespace ext::oneapi::experimental::ipc::event { +__SYCL_EXPORT handle get(const sycl::event &Evt); +__SYCL_EXPORT void put(handle &IpcHandle, const sycl::context &Ctx); +} // namespace ext::oneapi::experimental::ipc::event + namespace ext::oneapi::experimental::ipc { using handle_data_t = std::vector; @@ -71,6 +98,9 @@ struct handle { friend __SYCL_EXPORT handle physical_memory::get(physical_mem &physmem); friend __SYCL_EXPORT void physical_memory::put(handle &ipc_handle, const sycl::context &ctx); + friend __SYCL_EXPORT handle event::get(const sycl::event &Evt); + friend __SYCL_EXPORT void event::put(handle &IpcHandle, + const sycl::context &Ctx); }; } // namespace ext::oneapi::experimental::ipc diff --git a/sycl/include/sycl/ext/oneapi/experimental/ipc_event.hpp b/sycl/include/sycl/ext/oneapi/experimental/ipc_event.hpp new file mode 100644 index 0000000000000..f15e14b3738c4 --- /dev/null +++ b/sycl/include/sycl/ext/oneapi/experimental/ipc_event.hpp @@ -0,0 +1,114 @@ +//==------- ipc_event.hpp -- SYCL inter-process for events -----------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0) + +#include +#include +#include +#include +#include +#include +#include + +#include "detail/ipc_common.hpp" + +#include + +#if __has_include() +#include +#endif + +namespace sycl { +inline namespace _V1 { + +namespace detail { +// Bridge from the user-facing ipc::event::open overloads to a single +// __SYCL_EXPORT'd implementation. +__SYCL_EXPORT sycl::event openIPCEventHandle(const std::byte *HandleData, + size_t HandleDataSize, + const sycl::context &Ctx); + +// Producer-side IPC event creation. +__SYCL_EXPORT sycl::event makeIPCEvent(const sycl::context &Ctx); +} // namespace detail + +namespace ext::oneapi::experimental { + +// enable_ipc_key is defined in detail/ipc_common.hpp and shared with +// physical_mem; tag it as valid on sycl::event too. +template <> +struct is_property_key_of : std::true_type {}; + +// Minimum-viable make_event overload that handles only the enable_ipc +// property. The full reusable-events make_event lands with PR #22186; this +// overload will fold into it when that PR is merged. +template +inline sycl::event make_event(const sycl::context &Ctx, + PropertyListT Props = empty_properties_t{}) { + static_assert(is_property_list_v, + "Props must be a sycl::ext::oneapi::experimental::properties"); + (void)Props; + + if constexpr (PropertyListT::template has_property()) { + return sycl::detail::makeIPCEvent(Ctx); + } else { + static_assert( + sizeof(PropertyListT) == 0, + "make_event without the enable_ipc property is not yet supported. " + "The full reusable-events make_event lands with PR #22186."); + return sycl::event{}; + } +} + +} // namespace ext::oneapi::experimental + +namespace ext::oneapi::experimental::ipc::event { + +// get / put are exported from libsycl; declarations live in +// detail/ipc_common.hpp because they're friended on ipc::handle. + +inline void put(handle &IpcHandle) { + sycl::device Dev; + sycl::context Ctx = Dev.get_platform().khr_get_default_context(); + return put(IpcHandle, Ctx); +} + +inline sycl::event open(const ipc::handle_data_t &HandleData, + const sycl::context &Ctx) { + return sycl::detail::openIPCEventHandle(HandleData.data(), HandleData.size(), + Ctx); +} + +inline sycl::event open(const ipc::handle_data_t &HandleData) { + sycl::device Dev; + sycl::context Ctx = Dev.get_platform().khr_get_default_context(); + return open(HandleData, Ctx); +} + +#if __cpp_lib_span +inline sycl::event open(const ipc::handle_data_view_t &HandleDataView, + const sycl::context &Ctx) { + return sycl::detail::openIPCEventHandle(HandleDataView.data(), + HandleDataView.size(), Ctx); +} + +inline sycl::event open(const ipc::handle_data_view_t &HandleDataView) { + sycl::device Dev; + sycl::context Ctx = Dev.get_platform().khr_get_default_context(); + return open(HandleDataView, Ctx); +} +#endif + +} // namespace ext::oneapi::experimental::ipc::event +} // namespace _V1 +} // namespace sycl + +#endif diff --git a/sycl/include/sycl/ext/oneapi/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp index edb4e79be86a5..b426bdf119ec8 100644 --- a/sycl/include/sycl/ext/oneapi/properties/property.hpp +++ b/sycl/include/sycl/ext/oneapi/properties/property.hpp @@ -192,7 +192,7 @@ enum PropKind : uint32_t { MaximumSize = 47, ZeroInit = 48, FastLink = 49, - PhysicalMemoryEnableIPC = 50, + EnableIPC = 50, RegisterHostMemoryReadOnly = 51, // PropKindSize must always be the last value. PropKindSize = 52, diff --git a/sycl/include/sycl/ext/oneapi/virtual_mem/physical_mem.hpp b/sycl/include/sycl/ext/oneapi/virtual_mem/physical_mem.hpp index 3af65f44e8207..0454ba8c3aea4 100644 --- a/sycl/include/sycl/ext/oneapi/virtual_mem/physical_mem.hpp +++ b/sycl/include/sycl/ext/oneapi/virtual_mem/physical_mem.hpp @@ -13,6 +13,7 @@ #include #include #include +#include #include namespace sycl { @@ -26,13 +27,6 @@ namespace ext::oneapi::experimental { enum class address_access_mode : char { none = 0, read = 1, read_write = 2 }; -struct enable_ipc_key : detail::compile_time_property_key< - detail::PropKind::PhysicalMemoryEnableIPC> { - using value_t = property_value; -}; - -inline constexpr enable_ipc_key::value_t enable_ipc; - class __SYCL_EXPORT physical_mem : public sycl::detail::OwnerLessBase { friend sycl::detail::ImplUtils; diff --git a/sycl/include/sycl/info/aspects.def b/sycl/include/sycl/info/aspects.def index 960493719e83f..2b998f10144ea 100644 --- a/sycl/include/sycl/info/aspects.def +++ b/sycl/include/sycl/info/aspects.def @@ -92,3 +92,4 @@ __SYCL_ASPECT(ext_intel_eus_per_xe_core, 94) __SYCL_ASPECT(ext_intel_max_lanes_per_hw_thread, 95) __SYCL_ASPECT(ext_oneapi_ipc_physical_memory, 96) __SYCL_ASPECT(ext_oneapi_register_host_memory, 97) +__SYCL_ASPECT(ext_oneapi_ipc_event, 98) diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt index b7c119e177296..505ce5756d07d 100644 --- a/sycl/source/CMakeLists.txt +++ b/sycl/source/CMakeLists.txt @@ -359,6 +359,7 @@ set(SYCL_COMMON_SOURCES "handler.cpp" "image.cpp" "interop_handle.cpp" + "ipc_event.cpp" "ipc_memory.cpp" "ipc_physical_memory.cpp" "kernel.cpp" diff --git a/sycl/source/detail/device_impl.hpp b/sycl/source/detail/device_impl.hpp index b92710c6d226d..214ed3482bf9d 100644 --- a/sycl/source/detail/device_impl.hpp +++ b/sycl/source/detail/device_impl.hpp @@ -1622,6 +1622,10 @@ class device_impl { return get_info_impl_nocheck() .value_or(0); } + CASE(ext_oneapi_ipc_event) { + return get_info_impl_nocheck() + .value_or(0); + } CASE(ext_oneapi_device_wait) { return get_info_impl_nocheck() .value_or(0); diff --git a/sycl/source/detail/event_impl.cpp b/sycl/source/detail/event_impl.cpp index 6b7957bf781bd..e2dedcfed5051 100644 --- a/sycl/source/detail/event_impl.cpp +++ b/sycl/source/detail/event_impl.cpp @@ -167,6 +167,16 @@ event_impl::event_impl(ur_event_handle_t Event, const context &SyclContext, } } +event_impl::event_impl(ur_event_handle_t Event, const context &SyclContext, + IPCRole Role, private_tag tag) + : event_impl(Event, SyclContext, tag) { + // ext_oneapi_ipc_enabled() is the predicate for "can be exported via + // ipc::event::get()"; the spec restricts that to producer-side events + // created with the enable_ipc property. + MIPCEnabled = (Role == IPCRole::Producer); + MOpenedFromIpc = (Role == IPCRole::Imported); +} + event_impl::event_impl(queue_impl &Queue, private_tag) : MQueue{Queue.weak_from_this()}, MIsProfilingEnabled{Queue.MIsProfilingEnabled} { diff --git a/sycl/source/detail/event_impl.hpp b/sycl/source/detail/event_impl.hpp index 4e163660d7e20..31347495ec71e 100644 --- a/sycl/source/detail/event_impl.hpp +++ b/sycl/source/detail/event_impl.hpp @@ -69,6 +69,17 @@ class event_impl { /// \param SyclContext is an instance of SYCL context. event_impl(ur_event_handle_t Event, const context &SyclContext, private_tag); + /// IPC role of an event constructed from a UR handle. + enum class IPCRole { + /// Producer-side event created with enable_ipc. + Producer, + /// Consumer-side event imported via ipc::event::open(). + Imported, + }; + + event_impl(ur_event_handle_t Event, const context &SyclContext, IPCRole Role, + private_tag); + event_impl(queue_impl &Queue, private_tag); event_impl(HostEventState State, private_tag); @@ -96,6 +107,25 @@ class event_impl { private_tag{}); } + /// Producer-side IPC event factory: adopts a UR handle that backs + /// make_event(..., enable_ipc). Reports ext_oneapi_ipc_enabled() == true. + static std::shared_ptr + create_ipc_producer_event(ur_event_handle_t Event, + const context &SyclContext) { + return std::make_shared(Event, SyclContext, IPCRole::Producer, + private_tag{}); + } + + /// Consumer-side IPC event factory: adopts a UR handle returned by + /// ipc::event::open(). Imported events cannot be re-exported, so they + /// report ext_oneapi_ipc_enabled() == false. + static std::shared_ptr + create_ipc_imported_event(ur_event_handle_t Event, + const context &SyclContext) { + return std::make_shared(Event, SyclContext, IPCRole::Imported, + private_tag{}); + } + /// Sets a queue associated with the event /// /// Please note that this function changes the event state @@ -407,6 +437,25 @@ class event_impl { // storage. std::vector> MWeakPostCompleteEvents; + /// Both IPC flags are written exactly once, by the constructor that + /// the create_ipc_*_event factories invoke. After the + /// shared_ptr is returned to the caller these fields are + /// observed only read-only, so concurrent reads from any number of + /// threads are safe without synchronisation: the construction-then- + /// publish ordering is established by std::make_shared's release on + /// the produced shared_ptr. + + /// Backs event::ext_oneapi_ipc_enabled(). + bool MIPCEnabled = false; + + /// True only for events imported via ipc::event::open(). + bool MOpenedFromIpc = false; + +public: + bool isIPCEnabled() const noexcept { return MIPCEnabled; } + bool isOpenedFromIpc() const noexcept { return MOpenedFromIpc; } + +protected: /// Indicates that the task associated with this event has been submitted by /// the queue to the device. std::atomic MIsFlushed = false; diff --git a/sycl/source/detail/ur_device_info_ret_types.inc b/sycl/source/detail/ur_device_info_ret_types.inc index 5dd0829a2584e..a47ef3ddba7ca 100644 --- a/sycl/source/detail/ur_device_info_ret_types.inc +++ b/sycl/source/detail/ur_device_info_ret_types.inc @@ -172,6 +172,7 @@ MAP(UR_DEVICE_INFO_MAX_LANES_PER_HW_THREAD, uint32_t) // instead. MAP(UR_DEVICE_INFO_2D_BLOCK_ARRAY_CAPABILITIES_EXP, ur_exp_device_2d_block_array_capability_flags_t) MAP(UR_DEVICE_INFO_IPC_MEMORY_SUPPORT_EXP, ur_bool_t) +MAP(UR_DEVICE_INFO_IPC_EVENT_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/event.cpp b/sycl/source/event.cpp index d7e4389600a7d..4d4555682832e 100644 --- a/sycl/source/event.cpp +++ b/sycl/source/event.cpp @@ -114,6 +114,10 @@ backend event::get_backend() const noexcept try { std::abort(); } +bool event::ext_oneapi_ipc_enabled() const noexcept { + return impl->isIPCEnabled(); +} + ur_native_handle_t event::getNative() const { return impl->getNative(); } std::vector event::getNativeVector() const { diff --git a/sycl/source/ipc_event.cpp b/sycl/source/ipc_event.cpp new file mode 100644 index 0000000000000..fdba9ae537923 --- /dev/null +++ b/sycl/source/ipc_event.cpp @@ -0,0 +1,131 @@ +//==------- ipc_event.cpp -- SYCL inter-process for events -----------------==// +// +// 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 + +#include +#include + +namespace sycl { +inline namespace _V1 { + +namespace detail { + +// All devices in the context must support the aspect; spec: +// "_Throws:_ ... `errc::feature_not_supported` ... if not all devices that +// are part of `ctx` have `aspect::ext_oneapi_ipc_event`." +static void requireIPCEventAspect(const sycl::context &Ctx) { + for (const sycl::device &Dev : Ctx.get_devices()) { + if (!Dev.has(aspect::ext_oneapi_ipc_event)) + throw sycl::exception(sycl::make_error_code(errc::feature_not_supported), + "Not all devices in the context support " + "aspect::ext_oneapi_ipc_event."); + } +} + +__SYCL_EXPORT sycl::event openIPCEventHandle(const std::byte *HandleData, + size_t HandleDataSize, + const sycl::context &Ctx) { + requireIPCEventAspect(Ctx); + + auto CtxImpl = sycl::detail::getSyclObjImpl(Ctx); + sycl::detail::adapter_impl &Adapter = CtxImpl->getAdapter(); + + ur_event_handle_t UrEvent = nullptr; + ur_result_t UrRes = + Adapter.call_nocheck( + CtxImpl->getHandleRef(), HandleData, HandleDataSize, &UrEvent); + if (UrRes == UR_RESULT_ERROR_INVALID_VALUE) + throw sycl::exception( + sycl::make_error_code(errc::invalid), + "HandleData size does not correspond to the target platform's " + "IPC event handle size."); + Adapter.checkUrResult(UrRes); + assert(UrEvent && "urIPCOpenEventHandleExp returned success with null event"); + + // event_impl adopts the UR handle; release it directly if that step + // throws so the import doesn't leak. + try { + auto EventImpl = + sycl::detail::event_impl::create_ipc_imported_event(UrEvent, Ctx); + return sycl::detail::createSyclObjFromImpl(EventImpl); + } catch (...) { + Adapter.call_nocheck(UrEvent); + throw; + } +} + +__SYCL_EXPORT sycl::event makeIPCEvent(const sycl::context &Ctx) { + requireIPCEventAspect(Ctx); + + auto CtxImpl = sycl::detail::getSyclObjImpl(Ctx); + sycl::detail::adapter_impl &Adapter = CtxImpl->getAdapter(); + + // UR's urEventCreateExp takes a device parameter; pick the first device + // in the context. All devices in the context have the IPC event aspect + // (validated above), so any choice is valid. + const std::vector Devs = Ctx.get_devices(); + assert(!Devs.empty() && "context has no devices"); + + ur_exp_event_desc_t Desc{UR_STRUCTURE_TYPE_EXP_EVENT_DESC, nullptr, + UR_EXP_EVENT_FLAG_IPC_EXP}; + + ur_event_handle_t UrEvent = nullptr; + Adapter.call( + CtxImpl->getHandleRef(), getSyclObjImpl(Devs.front())->getHandleRef(), + &Desc, &UrEvent); + assert(UrEvent && "urEventCreateExp returned success with null event"); + + try { + auto EventImpl = + sycl::detail::event_impl::create_ipc_producer_event(UrEvent, Ctx); + return sycl::detail::createSyclObjFromImpl(EventImpl); + } catch (...) { + Adapter.call_nocheck(UrEvent); + throw; + } +} + +} // namespace detail + +namespace ext::oneapi::experimental::ipc::event { + +__SYCL_EXPORT handle get(const sycl::event &Evt) { + if (!Evt.ext_oneapi_ipc_enabled()) + throw sycl::exception( + sycl::make_error_code(errc::invalid), + "Event was not created with the enable_ipc property."); + + auto EvtImpl = sycl::detail::getSyclObjImpl(Evt); + sycl::detail::context_impl &CtxImpl = EvtImpl->getContextImpl(); + sycl::detail::adapter_impl &Adapter = CtxImpl.getAdapter(); + + void *HandlePtr = nullptr; + size_t HandleSize = 0; + Adapter.call( + EvtImpl->getHandle(), &HandlePtr, &HandleSize); + + return {HandlePtr, HandleSize}; +} + +__SYCL_EXPORT void put(handle &IpcHandle, const sycl::context &Ctx) { + auto CtxImpl = sycl::detail::getSyclObjImpl(Ctx); + sycl::detail::adapter_impl &Adapter = CtxImpl->getAdapter(); + + Adapter.call( + CtxImpl->getHandleRef(), IpcHandle.MData); +} + +} // namespace ext::oneapi::experimental::ipc::event +} // namespace _V1 +} // namespace sycl diff --git a/sycl/test-e2e/Experimental/Inputs/ipc_event_l0_signal.hpp b/sycl/test-e2e/Experimental/Inputs/ipc_event_l0_signal.hpp new file mode 100644 index 0000000000000..49e91a43e717c --- /dev/null +++ b/sycl/test-e2e/Experimental/Inputs/ipc_event_l0_signal.hpp @@ -0,0 +1,115 @@ +// Test-only helper that signals a SYCL IPC event through Level Zero interop by +// appending a barrier on an immediate command list. + +#pragma once + +#include +#include + +#include + +#include +#include +#include + +namespace ipc_event_test { + +// Create an in-order immediate command list on the device's compute engine. +inline ze_command_list_handle_t +createImmediateComputeCmdList(const sycl::context &Ctx, + const sycl::device &Dev) { + // Tests link the L0 loader statically; the driver must be initialised + // explicitly before any L0 entry points are called. + zeInit(ZE_INIT_FLAG_GPU_ONLY); + + ze_context_handle_t ZeCtx = + sycl::get_native(Ctx); + ze_device_handle_t ZeDev = + sycl::get_native(Dev); + + uint32_t NumGroups = 0; + if (zeDeviceGetCommandQueueGroupProperties(ZeDev, &NumGroups, nullptr) != + ZE_RESULT_SUCCESS) { + std::fprintf(stderr, "ipc_event_test: zeDeviceGetCommandQueueGroup* " + "failed\n"); + std::exit(2); + } + + std::vector Props( + NumGroups, {ZE_STRUCTURE_TYPE_COMMAND_QUEUE_GROUP_PROPERTIES}); + zeDeviceGetCommandQueueGroupProperties(ZeDev, &NumGroups, Props.data()); + + uint32_t ComputeOrdinal = 0; + for (uint32_t I = 0; I < NumGroups; ++I) { + if (Props[I].flags & ZE_COMMAND_QUEUE_GROUP_PROPERTY_FLAG_COMPUTE) { + ComputeOrdinal = I; + break; + } + } + + ze_command_queue_desc_t QDesc{ZE_STRUCTURE_TYPE_COMMAND_QUEUE_DESC}; + QDesc.ordinal = ComputeOrdinal; + QDesc.mode = ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS; + QDesc.flags = ZE_COMMAND_QUEUE_FLAG_IN_ORDER; + + ze_command_list_handle_t CmdList = nullptr; + if (zeCommandListCreateImmediate(ZeCtx, ZeDev, &QDesc, &CmdList) != + ZE_RESULT_SUCCESS) { + std::fprintf(stderr, "ipc_event_test: zeCommandListCreateImmediate " + "failed\n"); + std::exit(2); + } + return CmdList; +} + +// Signal SyclEvent by appending a barrier and then host-synchronize before +// returning (the signal has completed once this returns). +inline void signalEventViaLevelZero(sycl::event &SyclEvent, + const sycl::context &Ctx, + const sycl::device &Dev) { + ze_event_handle_t ZeEvt = + sycl::get_native(SyclEvent); + ze_command_list_handle_t CmdList = createImmediateComputeCmdList(Ctx, Dev); + + if (zeCommandListAppendBarrier(CmdList, ZeEvt, 0, nullptr) != + ZE_RESULT_SUCCESS) { + std::fprintf(stderr, "ipc_event_test: zeCommandListAppendBarrier " + "failed\n"); + std::exit(2); + } + if (zeEventHostSynchronize(ZeEvt, UINT64_MAX) != ZE_RESULT_SUCCESS) { + std::fprintf(stderr, "ipc_event_test: zeEventHostSynchronize failed\n"); + std::exit(2); + } + zeCommandListDestroy(CmdList); +} + +// Submit a barrier that signals SyclEvent only after DepEvent completes, then +// return WITHOUT host-synchronizing: SyclEvent may still be unsignaled when +// this returns. The signal is thus ordered strictly after DepEvent, so the +// only synchronization that guarantees DepEvent's effects to a remote waiter +// is a wait on SyclEvent itself. +// +// The returned command list must be kept alive until the signal has completed +// (e.g. after a later wait on SyclEvent) and then destroyed by the caller via +// zeCommandListDestroy. +inline ze_command_list_handle_t +submitSignalDependentOnEvent(sycl::event &SyclEvent, sycl::event &DepEvent, + const sycl::context &Ctx, + const sycl::device &Dev) { + ze_event_handle_t ZeEvt = + sycl::get_native(SyclEvent); + ze_event_handle_t ZeDep = + sycl::get_native(DepEvent); + ze_command_list_handle_t CmdList = createImmediateComputeCmdList(Ctx, Dev); + + if (zeCommandListAppendBarrier(CmdList, ZeEvt, 1, &ZeDep) != + ZE_RESULT_SUCCESS) { + std::fprintf(stderr, "ipc_event_test: zeCommandListAppendBarrier (with " + "dependency) failed\n"); + std::exit(2); + } + return CmdList; +} + +} // namespace ipc_event_test diff --git a/sycl/test-e2e/Experimental/Inputs/ipc_event_sentinel.hpp b/sycl/test-e2e/Experimental/Inputs/ipc_event_sentinel.hpp new file mode 100644 index 0000000000000..4b5ee47303d14 --- /dev/null +++ b/sycl/test-e2e/Experimental/Inputs/ipc_event_sentinel.hpp @@ -0,0 +1,73 @@ +// File-based rendezvous helpers shared by the cross-process IPC event tests. +// Paths are relative to the working directory; each test should prefix its +// file names to avoid collisions with other tests. + +#pragma once + +#include + +#include +#include +#include +#include +#include +#include +#include + +namespace ipc_event_test { + +namespace ipc = sycl::ext::oneapi::experimental::ipc; + +inline void waitForFile(const std::string &Path, int TimeoutSecs = 30) { + for (int i = 0; i < TimeoutSecs * 100; ++i) { + if (std::ifstream{Path}.good()) + return; + timespec Ts{0, 10'000'000}; // 10 ms + nanosleep(&Ts, nullptr); + } + std::fprintf(stderr, "Timeout waiting for %s\n", Path.c_str()); + std::exit(1); +} + +inline void touchFile(const std::string &Path) { + std::ofstream F(Path); + if (!F) + std::fprintf(stderr, "Failed to create %s\n", Path.c_str()); +} + +// Remove any leftover sentinel files from a previous run so that waitForFile +// does not match a stale file. Call this in the producer before it spawns the +// consumer; missing files are ignored. Paths are relative to the working +// directory, which lit does not clean between reruns. +inline void removeStaleFiles(std::initializer_list Paths) { + for (const char *P : Paths) + std::remove(P); +} + +// Layout: [size_t size][size bytes]. +inline void writeHandleFile(const std::string &Path, const ipc::handle &H) { + ipc::handle_data_t Bytes = H.data(); + size_t Sz = Bytes.size(); + std::ofstream FS(Path, std::ios::out | std::ios::binary); + FS.write(reinterpret_cast(&Sz), sizeof(Sz)); + FS.write(reinterpret_cast(Bytes.data()), Sz); + if (!FS) { + std::fprintf(stderr, "Failed to write handle to %s\n", Path.c_str()); + std::exit(1); + } +} + +inline ipc::handle_data_t readHandleFile(const std::string &Path) { + std::ifstream FS(Path, std::ios::in | std::ios::binary); + if (!FS) { + std::fprintf(stderr, "Failed to open handle file %s\n", Path.c_str()); + std::exit(1); + } + size_t Sz = 0; + FS.read(reinterpret_cast(&Sz), sizeof(Sz)); + std::unique_ptr Buf{new std::byte[Sz]}; + FS.read(reinterpret_cast(Buf.get()), Sz); + return ipc::handle_data_t(Buf.get(), Buf.get() + Sz); +} + +} // namespace ipc_event_test diff --git a/sycl/test-e2e/Experimental/ipc_event_basic.cpp b/sycl/test-e2e/Experimental/ipc_event_basic.cpp new file mode 100644 index 0000000000000..b761e7a66f345 --- /dev/null +++ b/sycl/test-e2e/Experimental/ipc_event_basic.cpp @@ -0,0 +1,42 @@ +// REQUIRES: aspect-ext_oneapi_ipc_event +// REQUIRES: level_zero_v2_adapter + +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +// Basic invariants of the SYCL IPC events surface that don't require +// signaling: aspect query, make_event creating an IPC-capable event, +// ext_oneapi_ipc_enabled() reporting correctly on regular vs IPC events. + +#include +#include + +namespace exp = sycl::ext::oneapi::experimental; + +int main() { + sycl::queue Q; + sycl::device Dev = Q.get_device(); + sycl::context Ctx = Q.get_context(); + + // 1. Aspect is true for this device. + if (!Dev.has(sycl::aspect::ext_oneapi_ipc_event)) + return 1; + + // 2. make_event(enable_ipc) produces an IPC-capable event. + sycl::event IpcEvt = exp::make_event(Ctx, exp::properties{exp::enable_ipc}); + if (!IpcEvt.ext_oneapi_ipc_enabled()) + return 2; + + // 3. A regular event is NOT IPC-capable. + sycl::event Plain = Q.submit([&](sycl::handler &H) { H.host_task([] {}); }); + Plain.wait(); + if (Plain.ext_oneapi_ipc_enabled()) + return 3; + + // 4. A default-constructed event is NOT IPC-capable. + sycl::event Default; + if (Default.ext_oneapi_ipc_enabled()) + return 4; + + return 0; +} diff --git a/sycl/test-e2e/Experimental/ipc_event_consumer_signal.cpp b/sycl/test-e2e/Experimental/ipc_event_consumer_signal.cpp new file mode 100644 index 0000000000000..0efa294750557 --- /dev/null +++ b/sycl/test-e2e/Experimental/ipc_event_consumer_signal.cpp @@ -0,0 +1,155 @@ +// REQUIRES: aspect-ext_oneapi_ipc_event +// REQUIRES: aspect-ext_oneapi_ipc_memory +// REQUIRES: level_zero_v2_adapter +// REQUIRES: arch-intel_gpu_bmg_g21 || arch-intel_gpu_bmg_g31 +// UNSUPPORTED: windows +// UNSUPPORTED-INTENDED: Cross-process IPC test relies on POSIX semantics. + +// RUN: %{build} -lze_loader %level_zero_options -o %t.out +// RUN: %{run} %t.out + +// The consumer signals an IPC event and the producer waits on it, with the +// result verified through a shared buffer. +// +// The producer owns the event and an IPC-shared USM buffer. The consumer opens +// both, enqueues a kernel that writes a sentinel into the buffer, and signals +// the event ordered after that kernel without waiting for it. The producer +// waits on the event and reads the buffer back. The event is the only thing +// ordering the write before the read: the "signal submitted" file does not +// imply completion. The buffer is poisoned first, so a wait that returns +// before the work completes reads the wrong value instead of hanging. +// +// clang-format off +// Sentinel protocol: +// sigdata_handles_ready producer -> consumer handles written, safe to open +// sigdata_signal_done consumer -> producer signalling command submitted +// sigdata_producer_synced producer -> consumer producer's wait done +// sigdata_consumer_done consumer -> producer consumer exited cleanly +// clang-format on + +#include "Inputs/ipc_event_l0_signal.hpp" +#include "Inputs/ipc_event_sentinel.hpp" +#include +#include +#include +#include + +#include +#include +#include + +#if defined(__linux__) +#include +#include +#endif + +namespace exp = sycl::ext::oneapi::experimental; +namespace ipc = sycl::ext::oneapi::experimental::ipc; + +// Number of int elements in the shared buffer and the value the consumer +// writes into every element. The poison value pre-fill must differ from it. +static constexpr size_t NumElems = 64; +static constexpr int Sentinel = 0x5A5A5A5A; +static constexpr int Poison = 0x0BADCAFE; + +static int producer(const std::string &Exe) { + sycl::queue Q; + sycl::device Dev = Q.get_device(); + sycl::context Ctx = Q.get_context(); + +#if defined(__linux__) + // Allow the unrelated consumer to pidfd_getfd into this process. + prctl(PR_SET_PTRACER, PR_SET_PTRACER_ANY); +#endif + + // lit does not clean the working directory between reruns; remove any stale + // sentinels so waitForFile does not match a leftover file from a prior run. + ipc_event_test::removeStaleFiles( + {"sigdata_handles_ready", "sigdata_signal_done", + "sigdata_producer_synced", "sigdata_consumer_done"}); + + sycl::event Evt = exp::make_event(Ctx, exp::properties{exp::enable_ipc}); + ipc::handle EvtHandle = ipc::event::get(Evt); + + // Allocate and poison the shared buffer before exporting it. + int *Buf = sycl::malloc_device(NumElems, Q); + Q.fill(Buf, Poison, NumElems).wait(); + ipc::handle MemHandle = ipc::memory::get(Buf, Ctx); + + ipc_event_test::writeHandleFile("sigdata_event.bin", EvtHandle); + ipc_event_test::writeHandleFile("sigdata_mem.bin", MemHandle); + ipc_event_test::touchFile("sigdata_handles_ready"); + + std::system((Exe + " consumer &").c_str()); + + ipc_event_test::waitForFile("sigdata_signal_done"); + Evt.wait(); + + // The wait returned, so the signal (ordered after the consumer's kernel) + // has completed and the sentinel must be visible now. + std::vector Host(NumElems); + Q.memcpy(Host.data(), Buf, NumElems * sizeof(int)).wait(); + + int Rc = 0; + for (size_t I = 0; I < NumElems; ++I) { + if (Host[I] != Sentinel) { + std::cerr << "FAILED: Buf[" << I << "] = " << std::hex << Host[I] + << ", expected " << Sentinel << std::dec << "\n"; + Rc = 1; + break; + } + } + + ipc_event_test::touchFile("sigdata_producer_synced"); + ipc_event_test::waitForFile("sigdata_consumer_done"); + + ipc::memory::put(MemHandle, Ctx); + ipc::event::put(EvtHandle, Ctx); + sycl::free(Buf, Q); + + if (Rc == 0) + std::cout << "PASSED: consumer signal carried data\n"; + return Rc; +} + +static int consumer() { + sycl::queue Q; + sycl::device Dev = Q.get_device(); + sycl::context Ctx = Q.get_context(); + + ipc_event_test::waitForFile("sigdata_handles_ready"); + auto EvtBytes = ipc_event_test::readHandleFile("sigdata_event.bin"); + auto MemBytes = ipc_event_test::readHandleFile("sigdata_mem.bin"); + + sycl::event Imported = ipc::event::open(EvtBytes, Ctx); + void *Shared = ipc::memory::open(MemBytes, Ctx, Dev); + int *Buf = static_cast(Shared); + + // Enqueue the sentinel write but DO NOT host-synchronize it: the signal must + // be the only thing that orders this write before the producer's read. + sycl::event WriteEvt = Q.parallel_for( + sycl::range<1>(NumElems), [=](sycl::id<1> I) { Buf[I] = Sentinel; }); + + // Submit a signal on the imported event that depends on the write kernel, + // without waiting for it to complete. + ze_command_list_handle_t SignalList = + ipc_event_test::submitSignalDependentOnEvent(Imported, WriteEvt, Ctx, + Dev); + ipc_event_test::touchFile("sigdata_signal_done"); + + // Stay alive while the producer's wait runs. + ipc_event_test::waitForFile("sigdata_producer_synced"); + + // The producer's wait has completed, so the signal (and thus the write) has + // completed; it is now safe to tear down the command list. + zeCommandListDestroy(SignalList); + ipc::memory::close(Shared, Ctx); + ipc_event_test::touchFile("sigdata_consumer_done"); + return 0; +} + +int main(int argc, char *argv[]) { + if (argc >= 2 && std::string(argv[1]) == "consumer") + return consumer(); + return producer(argv[0]); +} diff --git a/sycl/test-e2e/Experimental/ipc_event_lifetime.cpp b/sycl/test-e2e/Experimental/ipc_event_lifetime.cpp new file mode 100644 index 0000000000000..6b44465028b81 --- /dev/null +++ b/sycl/test-e2e/Experimental/ipc_event_lifetime.cpp @@ -0,0 +1,64 @@ +// REQUIRES: aspect-ext_oneapi_ipc_event +// REQUIRES: level_zero_v2_adapter + +// RUN: %{build} %level_zero_options -lze_loader -o %t.out +// RUN: %{run} %t.out + +// Lifetime invariants of the IPC events surface: +// * put() on the producer-side handle does not invalidate the producer +// event; a fresh get() on the same producer still yields a usable handle. +// * An imported event remains usable after the producer-side event is +// destroyed, as long as the imported event is still alive. + +#include "Inputs/ipc_event_l0_signal.hpp" +#include +#include + +#include + +namespace exp = sycl::ext::oneapi::experimental; +namespace ipc = sycl::ext::oneapi::experimental::ipc; + +int main() { + sycl::queue Q; + sycl::device Dev = Q.get_device(); + sycl::context Ctx = Q.get_context(); + + // 1. put does not invalidate the producer event. + { + sycl::event ProducerEvt = + exp::make_event(Ctx, exp::properties{exp::enable_ipc}); + ipc_event_test::signalEventViaLevelZero(ProducerEvt, Ctx, Dev); + + ipc::handle H1 = ipc::event::get(ProducerEvt); + ipc::event::put(H1, Ctx); + + // Producer is still usable: fresh get must succeed. + ipc::handle H2 = ipc::event::get(ProducerEvt); + if (H2.data().empty()) + return 1; + + sycl::event Imp = ipc::event::open(H2.data(), Ctx); + Imp.wait(); + ipc::event::put(H2, Ctx); + } + + // 2. Imported event survives release of the producer. + { + // Wrap the producer event in unique_ptr so we can drop it explicitly. + auto Producer = std::make_unique( + exp::make_event(Ctx, exp::properties{exp::enable_ipc})); + ipc_event_test::signalEventViaLevelZero(*Producer, Ctx, Dev); + + ipc::handle H = ipc::event::get(*Producer); + sycl::event Imported = ipc::event::open(H.data(), Ctx); + + // Drop the producer event; the imported event must remain usable. + Producer.reset(); + + Imported.wait(); + ipc::event::put(H, Ctx); + } + + return 0; +} diff --git a/sycl/test-e2e/Experimental/ipc_event_multi_process.cpp b/sycl/test-e2e/Experimental/ipc_event_multi_process.cpp new file mode 100644 index 0000000000000..58037440487cd --- /dev/null +++ b/sycl/test-e2e/Experimental/ipc_event_multi_process.cpp @@ -0,0 +1,117 @@ +// REQUIRES: aspect-ext_oneapi_ipc_event +// REQUIRES: level_zero_v2_adapter +// REQUIRES: arch-intel_gpu_bmg_g21 || arch-intel_gpu_bmg_g31 +// UNSUPPORTED: windows +// UNSUPPORTED-INTENDED: This test relies on POSIX +// semantics (std::system + binary self-exec). + +// RUN: %{build} -lze_loader %level_zero_options -o %t.out +// RUN: %{run} %t.out + +// Cross-process IPC event round trip: producer spawns a consumer with +// std::system() and transports the IPC handle bytes through a file. + +#include "Inputs/ipc_event_l0_signal.hpp" +#include +#include + +#include +#include +#include +#include +#include +#include +#include + +#if defined(__linux__) +#include +#include +#include +#endif // defined(__linux__) + +namespace exp = sycl::ext::oneapi::experimental; +namespace ipc = sycl::ext::oneapi::experimental::ipc; + +constexpr const char *CommsFile = "ipc_event_comms.bin"; + +int spawner(int argc, char *argv[]) { + assert(argc == 1); + sycl::queue Q; + sycl::device Dev = Q.get_device(); + sycl::context Ctx = Q.get_context(); + +#if defined(__linux__) + // Allow the consumer to pidfd_getfd into this process. + if (Q.get_backend() == sycl::backend::ext_oneapi_level_zero && + prctl(PR_SET_PTRACER, PR_SET_PTRACER_ANY) == -1) { + std::cerr << "Failed to set ptracer permissions!\n"; + return 1; + } +#endif + + sycl::event ProducerEvt = + exp::make_event(Ctx, exp::properties{exp::enable_ipc}); + ipc_event_test::signalEventViaLevelZero(ProducerEvt, Ctx, Dev); + + ipc::handle Handle = ipc::event::get(ProducerEvt); + ipc::handle_data_t Bytes = Handle.data(); + + // Write the handle bytes; layout matches readHandleFile in + // Inputs/ipc_event_sentinel.hpp: [size_t size][size bytes]. + { + std::ofstream FS(CommsFile, std::ios::out | std::ios::binary); + size_t HandleSize = Bytes.size(); + FS.write(reinterpret_cast(&HandleSize), sizeof(size_t)); + FS.write(reinterpret_cast(Bytes.data()), HandleSize); + } + + std::string Cmd = std::string{argv[0]} + " consumer"; + std::cout << "Spawning: " << Cmd << "\n"; + int Rc = std::system(Cmd.c_str()); + + ipc::event::put(Handle, Ctx); + std::remove(CommsFile); + + if (!WIFEXITED(Rc)) { + std::cerr << "Consumer process terminated abnormally, raw status " << Rc + << "\n"; + return 2; + } + if (WEXITSTATUS(Rc) != 0) { + std::cerr << "Consumer process failed with exit code " << WEXITSTATUS(Rc) + << "\n"; + return 2; + } + return 0; +} + +int consumer() { + sycl::queue Q; + sycl::device Dev = Q.get_device(); + sycl::context Ctx = Q.get_context(); + + std::ifstream FS(CommsFile, std::ios::in | std::ios::binary); + if (!FS) { + std::cerr << "consumer: failed to open " << CommsFile << "\n"; + return 1; + } + size_t HandleSize = 0; + FS.read(reinterpret_cast(&HandleSize), sizeof(size_t)); + if (HandleSize == 0) { + std::cerr << "consumer: zero-sized handle\n"; + return 2; + } + std::unique_ptr HandleBuffer{new std::byte[HandleSize]}; + FS.read(reinterpret_cast(HandleBuffer.get()), HandleSize); + + ipc::handle_data_t HandleVec(HandleBuffer.get(), + HandleBuffer.get() + HandleSize); + + sycl::event Imported = ipc::event::open(HandleVec, Ctx); + Imported.wait(); + return 0; +} + +int main(int argc, char *argv[]) { + return argc == 1 ? spawner(argc, argv) : consumer(); +} diff --git a/sycl/test-e2e/Experimental/ipc_event_negative.cpp b/sycl/test-e2e/Experimental/ipc_event_negative.cpp new file mode 100644 index 0000000000000..54866f4d49535 --- /dev/null +++ b/sycl/test-e2e/Experimental/ipc_event_negative.cpp @@ -0,0 +1,74 @@ +// REQUIRES: aspect-ext_oneapi_ipc_event +// REQUIRES: level_zero_v2_adapter + +// RUN: %{build} -lze_loader %level_zero_options -o %t.out +// RUN: %{run} %t.out + +// Negative path coverage of the SYCL IPC events surface: +// * ipc::event::get on a non-IPC event throws errc::invalid. +// * ipc::event::get on an imported event throws errc::invalid +// (imported events cannot be re-exported). +// * ipc::event::open with a buffer of the wrong size throws +// errc::invalid. + +#include "Inputs/ipc_event_l0_signal.hpp" +#include +#include + +namespace exp = sycl::ext::oneapi::experimental; +namespace ipc = sycl::ext::oneapi::experimental::ipc; + +static int expectInvalid(const sycl::exception &E) { + return E.code() == sycl::make_error_code(sycl::errc::invalid) ? 0 : 1; +} + +int main() { + sycl::queue Q; + sycl::device Dev = Q.get_device(); + sycl::context Ctx = Q.get_context(); + + // 1. get on a non-IPC event -> errc::invalid. + { + sycl::event Plain = Q.submit([&](sycl::handler &H) { H.host_task([] {}); }); + Plain.wait(); + try { + (void)ipc::event::get(Plain); + return 1; + } catch (const sycl::exception &E) { + if (expectInvalid(E)) + return 2; + } + } + + // 2. get on an imported event -> errc::invalid (cannot be re-exported). + { + sycl::event Producer = + exp::make_event(Ctx, exp::properties{exp::enable_ipc}); + ipc_event_test::signalEventViaLevelZero(Producer, Ctx, Dev); + ipc::handle H = ipc::event::get(Producer); + sycl::event Imported = ipc::event::open(H.data(), Ctx); + try { + (void)ipc::event::get(Imported); + ipc::event::put(H, Ctx); + return 7; + } catch (const sycl::exception &E) { + ipc::event::put(H, Ctx); + if (expectInvalid(E)) + return 8; + } + } + + // 3. open with a wrong-size buffer -> errc::invalid. + { + ipc::handle_data_t Bogus(7, std::byte{0x42}); + try { + (void)ipc::event::open(Bogus, Ctx); + return 3; + } catch (const sycl::exception &E) { + if (expectInvalid(E)) + return 4; + } + } + + return 0; +} diff --git a/sycl/test-e2e/Experimental/ipc_event_repeated_signal.cpp b/sycl/test-e2e/Experimental/ipc_event_repeated_signal.cpp new file mode 100644 index 0000000000000..810ace0893851 --- /dev/null +++ b/sycl/test-e2e/Experimental/ipc_event_repeated_signal.cpp @@ -0,0 +1,195 @@ +// REQUIRES: aspect-ext_oneapi_ipc_event +// REQUIRES: aspect-ext_oneapi_ipc_memory +// REQUIRES: level_zero_v2_adapter +// REQUIRES: arch-intel_gpu_bmg_g21 || arch-intel_gpu_bmg_g31 +// UNSUPPORTED: windows +// UNSUPPORTED-INTENDED: Cross-process IPC test relies on POSIX semantics. + +// RUN: %{build} -lze_loader %level_zero_options -o %t.out +// RUN: %{run} %t.out + +// A handle is opened once and waited on across two separate producer signals +// without reopening, with each wait verified through a shared buffer. +// +// The producer owns the event and an IPC-shared USM buffer. For each round it +// writes a distinct value into the buffer with a kernel and signals the event +// ordered after that kernel without waiting for it. The consumer opens the +// event and the buffer once, then per round waits and reads the buffer back. +// The event is the only thing ordering each round's write before the read, so +// a wait that returns before the work completes reads the wrong value instead +// of hanging. The rounds are serialized with a handshake so the producer does +// not overwrite round #1's value before the consumer reads it. +// +// clang-format off +// Sentinel protocol: +// repeat_handles_ready producer -> consumer handles written, safe to open +// repeat_signal1_ready producer -> consumer signal #1 command submitted +// repeat_consumed1 consumer -> producer value #1 read, safe to reuse +// repeat_signal2_ready producer -> consumer signal #2 command submitted +// repeat_consumed2 consumer -> producer both rounds done (verdict ready) +// repeat_consumer_failed consumer -> producer a data check failed (only on error) +// repeat_producer_done producer -> consumer verdict read, consumer may close +// repeat_consumer_done consumer -> producer handles closed, safe to release +// clang-format on +// +// The consumer runs in the background so its exit code is not observable by the +// producer; it reports a failed data check by creating repeat_consumer_failed +// before signaling repeat_consumed2, which the producer checks for. + +#include "Inputs/ipc_event_l0_signal.hpp" +#include "Inputs/ipc_event_sentinel.hpp" +#include +#include +#include +#include + +#include +#include +#include + +#if defined(__linux__) +#include +#include +#endif + +namespace exp = sycl::ext::oneapi::experimental; +namespace ipc = sycl::ext::oneapi::experimental::ipc; + +static constexpr size_t NumElems = 64; +static constexpr int Value1 = 0x11111111; +static constexpr int Value2 = 0x22222222; +static constexpr int Poison = 0x0BADCAFE; + +static int producer(const std::string &Exe) { + sycl::queue Q; + sycl::device Dev = Q.get_device(); + sycl::context Ctx = Q.get_context(); + +#if defined(__linux__) + prctl(PR_SET_PTRACER, PR_SET_PTRACER_ANY); +#endif + + // lit does not clean the working directory between reruns; remove any stale + // sentinels so waitForFile does not match a leftover file from a prior run. + ipc_event_test::removeStaleFiles( + {"repeat_handles_ready", "repeat_signal1_ready", "repeat_consumed1", + "repeat_signal2_ready", "repeat_consumed2", "repeat_consumer_failed", + "repeat_producer_done", "repeat_consumer_done"}); + + sycl::event Evt = exp::make_event(Ctx, exp::properties{exp::enable_ipc}); + ipc::handle EvtHandle = ipc::event::get(Evt); + + int *Buf = sycl::malloc_device(NumElems, Q); + Q.fill(Buf, Poison, NumElems).wait(); + ipc::handle MemHandle = ipc::memory::get(Buf, Ctx); + + ipc_event_test::writeHandleFile("repeat_event.bin", EvtHandle); + ipc_event_test::writeHandleFile("repeat_mem.bin", MemHandle); + ipc_event_test::touchFile("repeat_handles_ready"); + + std::system((Exe + " consumer &").c_str()); + + // Round 1: write Value1, signal #1 depends on that write. + sycl::event Write1 = Q.parallel_for(sycl::range<1>(NumElems), + [=](sycl::id<1> I) { Buf[I] = Value1; }); + ze_command_list_handle_t SignalList1 = + ipc_event_test::submitSignalDependentOnEvent(Evt, Write1, Ctx, Dev); + ipc_event_test::touchFile("repeat_signal1_ready"); + + // Wait until the consumer has read Value1 before overwriting the buffer. + ipc_event_test::waitForFile("repeat_consumed1"); + + // Round 2: write Value2, signal #2 depends on that write. + sycl::event Write2 = Q.parallel_for(sycl::range<1>(NumElems), + [=](sycl::id<1> I) { Buf[I] = Value2; }); + ze_command_list_handle_t SignalList2 = + ipc_event_test::submitSignalDependentOnEvent(Evt, Write2, Ctx, Dev); + ipc_event_test::touchFile("repeat_signal2_ready"); + + ipc_event_test::waitForFile("repeat_consumed2"); + + // The consumer signals a failed data check out-of-band, since its exit code + // is not observable (it runs in the background). + bool ConsumerFailed = std::ifstream{"repeat_consumer_failed"}.good(); + + // Both signals have completed (the consumer read both values), so the + // command lists are safe to destroy. + zeCommandListDestroy(SignalList1); + zeCommandListDestroy(SignalList2); + + // Let the consumer close its imported handles before we release the backing + // event and allocation, otherwise its close races with our free. + ipc_event_test::touchFile("repeat_producer_done"); + ipc_event_test::waitForFile("repeat_consumer_done"); + + ipc::memory::put(MemHandle, Ctx); + ipc::event::put(EvtHandle, Ctx); + sycl::free(Buf, Q); + + if (ConsumerFailed) { + std::cerr << "FAILED: consumer reported a data-check failure\n"; + return 1; + } + std::cout << "PASSED: repeated signal carried data\n"; + return 0; +} + +static int consumer() { + sycl::queue Q; + sycl::device Dev = Q.get_device(); + sycl::context Ctx = Q.get_context(); + + ipc_event_test::waitForFile("repeat_handles_ready"); + auto EvtBytes = ipc_event_test::readHandleFile("repeat_event.bin"); + auto MemBytes = ipc_event_test::readHandleFile("repeat_mem.bin"); + + // Open once, wait twice. + sycl::event Imported = ipc::event::open(EvtBytes, Ctx); + void *Shared = ipc::memory::open(MemBytes, Ctx, Dev); + int *Buf = static_cast(Shared); + + auto checkValue = [&](int Expected, const char *Round) -> bool { + std::vector Host(NumElems); + Q.memcpy(Host.data(), Buf, NumElems * sizeof(int)).wait(); + for (size_t I = 0; I < NumElems; ++I) { + if (Host[I] != Expected) { + std::cerr << "FAILED (" << Round << "): Buf[" << I << "] = " << std::hex + << Host[I] << ", expected " << Expected << std::dec << "\n"; + return false; + } + } + return true; + }; + + int Rc = 0; + + ipc_event_test::waitForFile("repeat_signal1_ready"); + Imported.wait(); + if (!checkValue(Value1, "signal #1")) + Rc = 1; + ipc_event_test::touchFile("repeat_consumed1"); + + if (Rc == 0) { + ipc_event_test::waitForFile("repeat_signal2_ready"); + Imported.wait(); + if (!checkValue(Value2, "signal #2")) + Rc = 1; + } + + // Report a failed data check out-of-band before releasing the producer, so + // its verdict survives this background process's discarded exit code. + if (Rc != 0) + ipc_event_test::touchFile("repeat_consumer_failed"); + ipc_event_test::touchFile("repeat_consumed2"); + + ipc_event_test::waitForFile("repeat_producer_done"); + ipc::memory::close(Shared, Ctx); + ipc_event_test::touchFile("repeat_consumer_done"); + return Rc; +} + +int main(int argc, char *argv[]) { + if (argc >= 2 && std::string(argv[1]) == "consumer") + return consumer(); + return producer(argv[0]); +} diff --git a/sycl/test-e2e/Experimental/ipc_event_round_trip.cpp b/sycl/test-e2e/Experimental/ipc_event_round_trip.cpp new file mode 100644 index 0000000000000..3f4104f518d7d --- /dev/null +++ b/sycl/test-e2e/Experimental/ipc_event_round_trip.cpp @@ -0,0 +1,87 @@ +// REQUIRES: aspect-ext_oneapi_ipc_event +// REQUIRES: level_zero_v2_adapter + +// DEFINE: %{cpp20} = %if cl_options %{/clang:-std=c++20%} %else %{-std=c++20%} + +// RUN: %{build} %level_zero_options -lze_loader -o %t.out +// RUN: %{run} %t.out +// RUN: %{build} %level_zero_options -lze_loader -DUSE_VIEW %{cpp20} -o %t.view.out +// RUN: %{run} %t.view.out + +// Full same-process Get / Open / Wait / Put round trip on an IPC event, +// signalled via the Level Zero interop helper. Covers: +// * the std::vector handle_data_t overload of open() and +// (with -DUSE_VIEW) the C++20 std::span handle_data_view_t overload, +// * default-context overloads of put() and open(), +// * multiple opens of the same handle yielding distinct events, +// * imported event reports ext_oneapi_ipc_enabled() == false (it cannot +// be re-exported), +// * imported event waits successfully because producer was signalled. + +#include "Inputs/ipc_event_l0_signal.hpp" +#include +#include + +namespace exp = sycl::ext::oneapi::experimental; +namespace ipc = sycl::ext::oneapi::experimental::ipc; + +int main() { + sycl::queue Q; + sycl::device Dev = Q.get_device(); + sycl::context Ctx = Q.get_context(); + + // 1. Create + signal the producer event. + sycl::event ProducerEvt = + exp::make_event(Ctx, exp::properties{exp::enable_ipc}); + ipc_event_test::signalEventViaLevelZero(ProducerEvt, Ctx, Dev); + + // 2. Get the IPC handle. + ipc::handle Handle = ipc::event::get(ProducerEvt); + + // 3. Open in this process via the requested overload. + sycl::event Imported1, Imported2; +#ifdef USE_VIEW + ipc::handle_data_view_t View = Handle.data_view(); + if (View.empty()) + return 10; + Imported1 = ipc::event::open(View, Ctx); + Imported2 = ipc::event::open(View, Ctx); +#else + ipc::handle_data_t Bytes = Handle.data(); + if (Bytes.empty()) + return 10; + Imported1 = ipc::event::open(Bytes, Ctx); + Imported2 = ipc::event::open(Bytes, Ctx); +#endif + + // 4. Imported events cannot be re-exported, so ext_oneapi_ipc_enabled() + // is false. + if (Imported1.ext_oneapi_ipc_enabled()) + return 1; + if (Imported2.ext_oneapi_ipc_enabled()) + return 2; + + // 5. They are not the same SYCL handle, even though they came from the + // same producer-side IPC bytes (the spec explicitly notes this: + // "each call to open may return a unique event object"). + if (Imported1 == Imported2) + return 3; + + // 6. Wait completes (producer was signalled). + Imported1.wait(); + Imported2.wait(); + + // 7. Default-context put compiles and runs. + ipc::event::put(Handle, Ctx); + + // Optional: also exercise the default-context open + put overloads. +#ifndef USE_VIEW + ipc::handle Handle2 = ipc::event::get(ProducerEvt); + ipc::handle_data_t Bytes2 = Handle2.data(); + sycl::event ImportedDefaultCtx = ipc::event::open(Bytes2); + ImportedDefaultCtx.wait(); + ipc::event::put(Handle2); +#endif + + return 0; +} diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 93782a73ca6a8..44f6b89e917b8 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3083,6 +3083,8 @@ _ZN4sycl3_V13ext6oneapi12experimental28unmap_external_linear_memoryEPvRKNS0_6dev _ZN4sycl3_V13ext6oneapi12experimental30supports_importing_handle_typeENS3_24external_mem_handle_typeERKNS0_6deviceE _ZN4sycl3_V13ext6oneapi12experimental3ipc15physical_memory3getERNS3_12physical_memE _ZN4sycl3_V13ext6oneapi12experimental3ipc15physical_memory3putERNS4_6handleERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental3ipc5event3getERKNS0_5eventE +_ZN4sycl3_V13ext6oneapi12experimental3ipc5event3putERNS4_6handleERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental3ipc6memory3getEPvRKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental3ipc6memory3putERNS4_6handleERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental3ipc6memory5closeEPvRKNS0_7contextE @@ -3268,6 +3270,7 @@ _ZN4sycl3_V16detail12buffer_plainC2EmRKNS0_7contextESt10unique_ptrINS1_19SYCLMem _ZN4sycl3_V16detail12buffer_plainC2EmmRKNS0_13property_listESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteIS7_EE _ZN4sycl3_V16detail12compile_implERKNS0_13kernel_bundleILNS0_12bundle_stateE0EEERKSt6vectorINS0_6deviceESaIS8_EERKNS0_13property_listE _ZN4sycl3_V16detail12isOutOfRangeENS0_3vecIiLi4EEENS0_15addressing_modeENS0_5rangeILi3EEE +_ZN4sycl3_V16detail12makeIPCEventERKNS0_7contextE _ZN4sycl3_V16detail12make_contextEmRKSt8functionIFvNS0_14exception_listEEENS0_7backendEbRKSt6vectorINS0_6deviceESaISA_EE _ZN4sycl3_V16detail13HandlerAccess10preProcessERNS0_7handlerENS1_19type_erased_cgfo_tyE _ZN4sycl3_V16detail13HandlerAccess11postProcessERNS0_7handlerENS1_19type_erased_cgfo_tyE @@ -3305,6 +3308,7 @@ _ZN4sycl3_V16detail17reduComputeWGSizeEmmRm _ZN4sycl3_V16detail18RangeRoundingTraceEv _ZN4sycl3_V16detail18get_kernel_id_implENS1_11string_viewE _ZN4sycl3_V16detail18make_kernel_bundleEmRKNS0_7contextEbNS0_12bundle_stateENS0_7backendE +_ZN4sycl3_V16detail18openIPCEventHandleEPKSt4bytemRKNS0_7contextE _ZN4sycl3_V16detail19defaultAsyncHandlerENS0_14exception_listE _ZN4sycl3_V16detail19getDeviceKernelInfoERKNS1_27compile_time_kernel_info_v123CompileTimeKernelInfoTyE _ZN4sycl3_V16detail19getDeviceKernelInfoESt17basic_string_viewIcSt11char_traitsIcEE @@ -3696,6 +3700,7 @@ _ZNK4sycl3_V15event15getNativeVectorEv _ZNK4sycl3_V15event18get_profiling_infoINS0_4info15event_profiling11command_endEEENS0_6detail28is_event_profiling_info_descIT_E11return_typeEv _ZNK4sycl3_V15event18get_profiling_infoINS0_4info15event_profiling13command_startEEENS0_6detail28is_event_profiling_info_descIT_E11return_typeEv _ZNK4sycl3_V15event18get_profiling_infoINS0_4info15event_profiling14command_submitEEENS0_6detail28is_event_profiling_info_descIT_E11return_typeEv +_ZNK4sycl3_V15event22ext_oneapi_ipc_enabledEv _ZNK4sycl3_V15event8get_infoINS0_4info5event15reference_countEEENS0_6detail18is_event_info_descIT_E11return_typeEv _ZNK4sycl3_V15event8get_infoINS0_4info5event24command_execution_statusEEENS0_6detail18is_event_info_descIT_E11return_typeEv _ZNK4sycl3_V15event9getNativeEv diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index e683e58d8c089..03be5130a300c 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -3901,6 +3901,7 @@ ?ext_oneapi_has_kernel@kernel_bundle_plain@detail@_V1@sycl@@AEAA_NVstring_view@234@@Z ?ext_oneapi_has_kernel@kernel_bundle_plain@detail@_V1@sycl@@QEAA_NAEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@Z ?ext_oneapi_index_within_platform@device@_V1@sycl@@QEBA_KXZ +?ext_oneapi_ipc_enabled@event@_V1@sycl@@QEBA_NXZ ?ext_oneapi_memcpy2d_impl@handler@_V1@sycl@@AEAAXPEAX_KPEBX111@Z ?ext_oneapi_memset2d_impl@handler@_V1@sycl@@AEAAXPEAX_KH11@Z ?ext_oneapi_owner_before@?$OwnerLessBase@Vcontext@_V1@sycl@@@detail@_V1@sycl@@QEBA_NAEBV?$weak_object_base@Vcontext@_V1@sycl@@@2oneapi@ext@34@@Z @@ -3970,6 +3971,7 @@ ?frexp_impl@detail@_V1@sycl@@YANNPEAH@Z ?get@context@_V1@sycl@@QEBAPEAU_cl_context@@XZ ?get@device@_V1@sycl@@QEBAPEAU_cl_device_id@@XZ +?get@event@ipc@experimental@oneapi@ext@_V1@sycl@@YA?AUhandle@234567@AEBV167@@Z ?get@ipc_memory@experimental@oneapi@ext@_V1@sycl@@YA?AUhandle@123456@PEAXAEBVcontext@56@@Z ?get@kernel@_V1@sycl@@QEBAPEAU_cl_kernel@@XZ ?get@memory@ipc@experimental@oneapi@ext@_V1@sycl@@YA?AUhandle@234567@PEAXAEBVcontext@67@@Z @@ -4210,6 +4212,7 @@ ?link_impl@detail@_V1@sycl@@YA?AV?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@AEBV?$vector@V?$kernel_bundle@$00@_V1@sycl@@V?$allocator@V?$kernel_bundle@$00@_V1@sycl@@@std@@@5@AEBV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@5@AEBVproperty_list@23@@Z ?link_impl@detail@_V1@sycl@@YA?AV?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@PEBV?$kernel_bundle@$00@23@_KAEBV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@5@_N@Z ?makeDir@OSUtil@detail@_V1@sycl@@SAHPEBD@Z +?makeIPCEvent@detail@_V1@sycl@@YA?AVevent@23@AEBVcontext@23@@Z ?make_context@detail@_V1@sycl@@YA?AVcontext@23@_KAEBV?$function@$$A6AXVexception_list@_V1@sycl@@@Z@std@@W4backend@23@_NAEBV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@6@@Z ?make_device@detail@_V1@sycl@@YA?AVdevice@23@_KW4backend@23@@Z ?make_device@detail@level_zero@oneapi@ext@_V1@sycl@@YA?AVdevice@56@AEBVplatform@56@_K@Z @@ -4275,6 +4278,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 +?openIPCEventHandle@detail@_V1@sycl@@YA?AVevent@23@PEBW4byte@std@@_KAEBVcontext@23@@Z ?openIPCMemHandle@detail@_V1@sycl@@YAPEAXPEBW4byte@std@@_KAEBVcontext@23@AEBVdevice@23@@Z ?openIPCPhysicalMemHandle@detail@_V1@sycl@@YA?AVphysical_mem@experimental@oneapi@ext@23@PEBW4byte@std@@_KAEBVcontext@23@AEBVdevice@23@@Z ?parallel_for@handler@_V1@sycl@@QEAAXV?$range@$00@23@Vkernel@23@@Z @@ -4295,6 +4299,7 @@ ?prepare_for_device_copy@experimental@oneapi@ext@_V1@sycl@@YAXPEBX_KAEBVqueue@45@@Z ?print_graph@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@IEBAXVstring_view@267@_N@Z ?print_graph@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@QEBAXV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@_N@Z +?put@event@ipc@experimental@oneapi@ext@_V1@sycl@@YAXAEAUhandle@234567@AEBVcontext@67@@Z ?put@ipc_memory@experimental@oneapi@ext@_V1@sycl@@YAXAEAUhandle@123456@AEBVcontext@56@@Z ?put@memory@ipc@experimental@oneapi@ext@_V1@sycl@@YAXAEAUhandle@234567@AEBVcontext@67@@Z ?put@physical_memory@ipc@experimental@oneapi@ext@_V1@sycl@@YAXAEAUhandle@234567@AEBVcontext@67@@Z