From a5edca2dc0ac5b36ff3046ed3207c1d06a9b9d56 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Mon, 15 Sep 2025 03:41:41 -0700 Subject: [PATCH] [SYCL][BindlessImages] Fix storing result events for semaphores 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 --- 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 bac765cb8d1e5..b9eb9fdf3c1e5 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -3684,11 +3684,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 && @@ -3698,11 +3704,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); }