From 0c7a1d91ea3a4ad66464bef1a8f3f6700b500bbb Mon Sep 17 00:00:00 2001 From: Steffen Larsen Date: Fri, 12 Sep 2025 01:11:31 +0200 Subject: [PATCH 1/2] [SYCL][BindlessImages] Fix external semaphore dependencies and return events (#20040) This commit fixes an issue where bindless images semaphore operations (signal/wait) would neither use dependency events of the submission nor return the corresponding event from the backend operation. This commit fixes both of these issues. --------- Signed-off-by: Larsen, Steffen (cherry picked from commit b578d545dde02747cba524473d7d7635cf149a46) --- sycl/source/detail/scheduler/commands.cpp | 6 +- .../Extensions/BindlessImages/CMakeLists.txt | 3 + .../Extensions/BindlessImages/Semaphores.cpp | 161 ++++++++++++++++++ sycl/unittests/Extensions/CMakeLists.txt | 1 + 4 files changed, 169 insertions(+), 2 deletions(-) create mode 100644 sycl/unittests/Extensions/BindlessImages/CMakeLists.txt create mode 100644 sycl/unittests/Extensions/BindlessImages/Semaphores.cpp diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index e427d851148c1..7abaa3593f60b 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -3713,7 +3713,8 @@ ur_result_t ExecCGCommand::enqueueImpQueue() { return Adapter .call_nocheck( MQueue->getHandleRef(), SemWait->getExternalSemaphore(), - OptWaitValue.has_value(), WaitValue, 0, nullptr, nullptr); + OptWaitValue.has_value(), WaitValue, RawEvents.size(), + RawEvents.data(), Event); } case CGType::SemaphoreSignal: { assert(MQueue && @@ -3726,7 +3727,8 @@ ur_result_t ExecCGCommand::enqueueImpQueue() { return Adapter .call_nocheck( MQueue->getHandleRef(), SemSignal->getExternalSemaphore(), - OptSignalValue.has_value(), SignalValue, 0, nullptr, nullptr); + OptSignalValue.has_value(), SignalValue, RawEvents.size(), + RawEvents.data(), Event); } case CGType::AsyncAlloc: { // NO-OP. Async alloc calls adapter immediately in order to return a valid diff --git a/sycl/unittests/Extensions/BindlessImages/CMakeLists.txt b/sycl/unittests/Extensions/BindlessImages/CMakeLists.txt new file mode 100644 index 0000000000000..3745a8cec0fbc --- /dev/null +++ b/sycl/unittests/Extensions/BindlessImages/CMakeLists.txt @@ -0,0 +1,3 @@ +add_sycl_unittest(BindlessImagesExtensionTests OBJECT + Semaphores.cpp +) diff --git a/sycl/unittests/Extensions/BindlessImages/Semaphores.cpp b/sycl/unittests/Extensions/BindlessImages/Semaphores.cpp new file mode 100644 index 0000000000000..adc9a0186d35b --- /dev/null +++ b/sycl/unittests/Extensions/BindlessImages/Semaphores.cpp @@ -0,0 +1,161 @@ +#include + +#include + +#include +#include +#include +#include +#include + +namespace syclexp = sycl::ext::oneapi::experimental; + +constexpr uint64_t WaitValue = 42; +constexpr uint64_t SignalValue = 24; + +thread_local int urBindlessImagesWaitExternalSemaphoreExp_counter = 0; +thread_local bool urBindlessImagesWaitExternalSemaphoreExp_expectHasWaitValue = + false; +inline ur_result_t +urBindlessImagesWaitExternalSemaphoreExp_replace(void *pParams) { + ++urBindlessImagesWaitExternalSemaphoreExp_counter; + ur_bindless_images_wait_external_semaphore_exp_params_t Params = + *reinterpret_cast< + ur_bindless_images_wait_external_semaphore_exp_params_t *>(pParams); + EXPECT_EQ(*Params.phasWaitValue, + urBindlessImagesWaitExternalSemaphoreExp_expectHasWaitValue); + if (urBindlessImagesWaitExternalSemaphoreExp_expectHasWaitValue) { + EXPECT_EQ(*Params.pwaitValue, WaitValue); + } + EXPECT_EQ(*Params.pnumEventsInWaitList, uint32_t{0}); + EXPECT_EQ(*Params.pphEventWaitList, nullptr); + EXPECT_NE(*Params.pphEvent, nullptr); + return UR_RESULT_SUCCESS; +} + +thread_local int urBindlessImagesSignalExternalSemaphoreExp_counter = 0; +thread_local bool + urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = false; +thread_local uint32_t + urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 0; +inline ur_result_t +urBindlessImagesSignalExternalSemaphoreExp_replace(void *pParams) { + ++urBindlessImagesSignalExternalSemaphoreExp_counter; + ur_bindless_images_signal_external_semaphore_exp_params_t Params = + *reinterpret_cast< + ur_bindless_images_signal_external_semaphore_exp_params_t *>(pParams); + EXPECT_EQ(*Params.phasSignalValue, + urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue); + if (urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue) { + EXPECT_EQ(*Params.psignalValue, SignalValue); + } + EXPECT_EQ(*Params.pnumEventsInWaitList, + urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents); + if (urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents) { + EXPECT_NE(*Params.pphEventWaitList, nullptr); + } else { + EXPECT_EQ(*Params.pphEventWaitList, nullptr); + } + EXPECT_NE(*Params.pphEvent, nullptr); + return UR_RESULT_SUCCESS; +} + +TEST(BindlessImagesExtensionTests, ExternalSemaphoreWait) { + sycl::unittest::UrMock<> Mock; + mock::getCallbacks().set_replace_callback( + "urBindlessImagesWaitExternalSemaphoreExp", + &urBindlessImagesWaitExternalSemaphoreExp_replace); + urBindlessImagesWaitExternalSemaphoreExp_counter = 0; + + sycl::queue Q; + + // Create a dummy external semaphore and set the raw handle to some dummy. + // The mock implementation should never access the handle, so this is safe. + int DummyInt = 0; + syclexp::external_semaphore DummySemaphore{}; + DummySemaphore.raw_handle = + reinterpret_cast(&DummyInt); + + DummySemaphore.handle_type = + syclexp::external_semaphore_handle_type::opaque_fd; + + urBindlessImagesWaitExternalSemaphoreExp_expectHasWaitValue = false; + Q.ext_oneapi_wait_external_semaphore(DummySemaphore); + EXPECT_EQ(urBindlessImagesWaitExternalSemaphoreExp_counter, 1); + + DummySemaphore.handle_type = + syclexp::external_semaphore_handle_type::timeline_fd; + + urBindlessImagesWaitExternalSemaphoreExp_expectHasWaitValue = true; + Q.ext_oneapi_wait_external_semaphore(DummySemaphore, WaitValue); + EXPECT_EQ(urBindlessImagesWaitExternalSemaphoreExp_counter, 2); +} + +TEST(BindlessImagesExtensionTests, ExternalSemaphoreSignal) { + sycl::unittest::UrMock<> Mock; + mock::getCallbacks().set_replace_callback( + "urBindlessImagesSignalExternalSemaphoreExp", + &urBindlessImagesSignalExternalSemaphoreExp_replace); + urBindlessImagesSignalExternalSemaphoreExp_counter = 0; + + sycl::queue Q; + + // Create a dummy external semaphore and set the raw handle to some dummy. + // The mock implementation should never access the handle, so this is safe. + int DummyInt1 = 0, DummyInt2 = 0; + syclexp::external_semaphore DummySemaphore{}; + DummySemaphore.raw_handle = + reinterpret_cast(&DummyInt1); + + // We create dummy events with dummy UR handles to make the runtime think we + // pass actual device events. + auto DummyEventImpl1 = sycl::detail::event_impl::create_device_event( + *sycl::detail::getSyclObjImpl(Q)); + auto DummyEventImpl2 = sycl::detail::event_impl::create_device_event( + *sycl::detail::getSyclObjImpl(Q)); + DummyEventImpl1->setHandle(reinterpret_cast(&DummyInt1)); + DummyEventImpl2->setHandle(reinterpret_cast(&DummyInt2)); + sycl::event DummyEvent1 = + sycl::detail::createSyclObjFromImpl(DummyEventImpl1); + sycl::event DummyEvent2 = + sycl::detail::createSyclObjFromImpl(DummyEventImpl2); + std::vector DummyEventList{DummyEvent1, DummyEvent2}; + + DummySemaphore.handle_type = + syclexp::external_semaphore_handle_type::opaque_fd; + + urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = false; + urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 0; + Q.ext_oneapi_signal_external_semaphore(DummySemaphore); + EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 1); + + urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = false; + urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 1; + Q.ext_oneapi_signal_external_semaphore(DummySemaphore, DummyEvent1); + EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 2); + + urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = false; + urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 2; + Q.ext_oneapi_signal_external_semaphore(DummySemaphore, DummyEventList); + EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 3); + + DummySemaphore.handle_type = + syclexp::external_semaphore_handle_type::timeline_fd; + + urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = true; + urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 0; + Q.ext_oneapi_signal_external_semaphore(DummySemaphore, SignalValue); + EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 4); + + urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = true; + urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 1; + Q.ext_oneapi_signal_external_semaphore(DummySemaphore, SignalValue, + DummyEvent1); + EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 5); + + urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = true; + urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 2; + Q.ext_oneapi_signal_external_semaphore(DummySemaphore, SignalValue, + DummyEventList); + EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 6); +} diff --git a/sycl/unittests/Extensions/CMakeLists.txt b/sycl/unittests/Extensions/CMakeLists.txt index b82c9f798a94c..da81e86c23f38 100644 --- a/sycl/unittests/Extensions/CMakeLists.txt +++ b/sycl/unittests/Extensions/CMakeLists.txt @@ -24,6 +24,7 @@ add_sycl_unittest(ExtensionsTests OBJECT RootGroup.cpp ) +add_subdirectory(BindlessImages) add_subdirectory(CommandGraph) add_subdirectory(VirtualFunctions) add_subdirectory(VirtualMemory) From 0be37b3731d5a0a72a0d410b37abf250adfcc5b4 Mon Sep 17 00:00:00 2001 From: Steffen Larsen Date: Mon, 15 Sep 2025 19:09:40 +0200 Subject: [PATCH 2/2] [SYCL][BindlessImages] Fix storing result events for semaphores (#20080) https://github.com/intel/llvm/pull/20040 addressed an issue where semaphore operations would not pass and retrieve events from semaphore operations. However, the changes did not correctly store the result events. This commit addresses this. Signed-off-by: Larsen, Steffen (cherry picked from commit 68f3fdf41ca373e413c74da2949d807d3d7d777f) --- sycl/source/detail/scheduler/commands.cpp | 20 ++++++-- .../Extensions/BindlessImages/Semaphores.cpp | 50 +++++++++++++++---- 2 files changed, 56 insertions(+), 14 deletions(-) diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index 7abaa3593f60b..3a885d0875acd 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -3710,11 +3710,17 @@ ur_result_t ExecCGCommand::enqueueImpQueue() { auto OptWaitValue = SemWait->getWaitValue(); uint64_t WaitValue = OptWaitValue.has_value() ? OptWaitValue.value() : 0; - return Adapter - .call_nocheck( + if (auto Result = Adapter.call_nocheck< + UrApiKind::urBindlessImagesWaitExternalSemaphoreExp>( MQueue->getHandleRef(), SemWait->getExternalSemaphore(), OptWaitValue.has_value(), WaitValue, RawEvents.size(), RawEvents.data(), Event); + Result != UR_RESULT_SUCCESS) + return Result; + + SetEventHandleOrDiscard(); + + return UR_RESULT_SUCCESS; } case CGType::SemaphoreSignal: { assert(MQueue && @@ -3724,11 +3730,17 @@ ur_result_t ExecCGCommand::enqueueImpQueue() { auto OptSignalValue = SemSignal->getSignalValue(); uint64_t SignalValue = OptSignalValue.has_value() ? OptSignalValue.value() : 0; - return Adapter - .call_nocheck( + if (auto Result = Adapter.call_nocheck< + UrApiKind::urBindlessImagesSignalExternalSemaphoreExp>( MQueue->getHandleRef(), SemSignal->getExternalSemaphore(), OptSignalValue.has_value(), SignalValue, RawEvents.size(), RawEvents.data(), Event); + Result != UR_RESULT_SUCCESS) + return Result; + + SetEventHandleOrDiscard(); + + return UR_RESULT_SUCCESS; } case CGType::AsyncAlloc: { // NO-OP. Async alloc calls adapter immediately in order to return a valid diff --git a/sycl/unittests/Extensions/BindlessImages/Semaphores.cpp b/sycl/unittests/Extensions/BindlessImages/Semaphores.cpp index adc9a0186d35b..16f6aa2917920 100644 --- a/sycl/unittests/Extensions/BindlessImages/Semaphores.cpp +++ b/sycl/unittests/Extensions/BindlessImages/Semaphores.cpp @@ -16,6 +16,8 @@ constexpr uint64_t SignalValue = 24; thread_local int urBindlessImagesWaitExternalSemaphoreExp_counter = 0; thread_local bool urBindlessImagesWaitExternalSemaphoreExp_expectHasWaitValue = false; +thread_local ur_event_handle_t + urBindlessImagesWaitExternalSemaphoreExp_lastEvent = nullptr; inline ur_result_t urBindlessImagesWaitExternalSemaphoreExp_replace(void *pParams) { ++urBindlessImagesWaitExternalSemaphoreExp_counter; @@ -30,6 +32,11 @@ urBindlessImagesWaitExternalSemaphoreExp_replace(void *pParams) { EXPECT_EQ(*Params.pnumEventsInWaitList, uint32_t{0}); EXPECT_EQ(*Params.pphEventWaitList, nullptr); EXPECT_NE(*Params.pphEvent, nullptr); + if (*Params.pphEvent) { + urBindlessImagesWaitExternalSemaphoreExp_lastEvent = + mock::createDummyHandle(); + **Params.pphEvent = urBindlessImagesWaitExternalSemaphoreExp_lastEvent; + } return UR_RESULT_SUCCESS; } @@ -38,6 +45,8 @@ thread_local bool urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = false; thread_local uint32_t urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 0; +thread_local ur_event_handle_t + urBindlessImagesSignalExternalSemaphoreExp_lastEvent = nullptr; inline ur_result_t urBindlessImagesSignalExternalSemaphoreExp_replace(void *pParams) { ++urBindlessImagesSignalExternalSemaphoreExp_counter; @@ -57,6 +66,11 @@ urBindlessImagesSignalExternalSemaphoreExp_replace(void *pParams) { EXPECT_EQ(*Params.pphEventWaitList, nullptr); } EXPECT_NE(*Params.pphEvent, nullptr); + if (*Params.pphEvent) { + urBindlessImagesSignalExternalSemaphoreExp_lastEvent = + mock::createDummyHandle(); + **Params.pphEvent = urBindlessImagesSignalExternalSemaphoreExp_lastEvent; + } return UR_RESULT_SUCCESS; } @@ -80,15 +94,19 @@ TEST(BindlessImagesExtensionTests, ExternalSemaphoreWait) { syclexp::external_semaphore_handle_type::opaque_fd; urBindlessImagesWaitExternalSemaphoreExp_expectHasWaitValue = false; - Q.ext_oneapi_wait_external_semaphore(DummySemaphore); + sycl::event E = Q.ext_oneapi_wait_external_semaphore(DummySemaphore); EXPECT_EQ(urBindlessImagesWaitExternalSemaphoreExp_counter, 1); + EXPECT_EQ(sycl::detail::getSyclObjImpl(E)->getHandle(), + urBindlessImagesWaitExternalSemaphoreExp_lastEvent); DummySemaphore.handle_type = syclexp::external_semaphore_handle_type::timeline_fd; urBindlessImagesWaitExternalSemaphoreExp_expectHasWaitValue = true; - Q.ext_oneapi_wait_external_semaphore(DummySemaphore, WaitValue); + E = Q.ext_oneapi_wait_external_semaphore(DummySemaphore, WaitValue); EXPECT_EQ(urBindlessImagesWaitExternalSemaphoreExp_counter, 2); + EXPECT_EQ(sycl::detail::getSyclObjImpl(E)->getHandle(), + urBindlessImagesWaitExternalSemaphoreExp_lastEvent); } TEST(BindlessImagesExtensionTests, ExternalSemaphoreSignal) { @@ -126,36 +144,48 @@ TEST(BindlessImagesExtensionTests, ExternalSemaphoreSignal) { urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = false; urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 0; - Q.ext_oneapi_signal_external_semaphore(DummySemaphore); + sycl::event E = Q.ext_oneapi_signal_external_semaphore(DummySemaphore); EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 1); + EXPECT_EQ(sycl::detail::getSyclObjImpl(E)->getHandle(), + urBindlessImagesSignalExternalSemaphoreExp_lastEvent); urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = false; urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 1; - Q.ext_oneapi_signal_external_semaphore(DummySemaphore, DummyEvent1); + E = Q.ext_oneapi_signal_external_semaphore(DummySemaphore, DummyEvent1); EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 2); + EXPECT_EQ(sycl::detail::getSyclObjImpl(E)->getHandle(), + urBindlessImagesSignalExternalSemaphoreExp_lastEvent); urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = false; urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 2; - Q.ext_oneapi_signal_external_semaphore(DummySemaphore, DummyEventList); + E = Q.ext_oneapi_signal_external_semaphore(DummySemaphore, DummyEventList); EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 3); + EXPECT_EQ(sycl::detail::getSyclObjImpl(E)->getHandle(), + urBindlessImagesSignalExternalSemaphoreExp_lastEvent); DummySemaphore.handle_type = syclexp::external_semaphore_handle_type::timeline_fd; urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = true; urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 0; - Q.ext_oneapi_signal_external_semaphore(DummySemaphore, SignalValue); + E = Q.ext_oneapi_signal_external_semaphore(DummySemaphore, SignalValue); EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 4); + EXPECT_EQ(sycl::detail::getSyclObjImpl(E)->getHandle(), + urBindlessImagesSignalExternalSemaphoreExp_lastEvent); urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = true; urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 1; - Q.ext_oneapi_signal_external_semaphore(DummySemaphore, SignalValue, - DummyEvent1); + E = Q.ext_oneapi_signal_external_semaphore(DummySemaphore, SignalValue, + DummyEvent1); EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 5); + EXPECT_EQ(sycl::detail::getSyclObjImpl(E)->getHandle(), + urBindlessImagesSignalExternalSemaphoreExp_lastEvent); urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = true; urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 2; - Q.ext_oneapi_signal_external_semaphore(DummySemaphore, SignalValue, - DummyEventList); + E = Q.ext_oneapi_signal_external_semaphore(DummySemaphore, SignalValue, + DummyEventList); EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 6); + EXPECT_EQ(sycl::detail::getSyclObjImpl(E)->getHandle(), + urBindlessImagesSignalExternalSemaphoreExp_lastEvent); }