Skip to content

Commit b578d54

Browse files
steffenlarsensteffenlarsen
authored
[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 <[email protected]>
1 parent a20b51a commit b578d54

File tree

4 files changed

+169
-2
lines changed

4 files changed

+169
-2
lines changed

‎sycl/source/detail/scheduler/commands.cpp

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -3687,7 +3687,8 @@ ur_result_t ExecCGCommand::enqueueImpQueue() {
36873687
return Adapter
36883688
.call_nocheck<UrApiKind::urBindlessImagesWaitExternalSemaphoreExp>(
36893689
MQueue->getHandleRef(), SemWait->getExternalSemaphore(),
3690-
OptWaitValue.has_value(), WaitValue, 0, nullptr, nullptr);
3690+
OptWaitValue.has_value(), WaitValue, RawEvents.size(),
3691+
RawEvents.data(), Event);
36913692
}
36923693
case CGType::SemaphoreSignal: {
36933694
assert(MQueue &&
@@ -3700,7 +3701,8 @@ ur_result_t ExecCGCommand::enqueueImpQueue() {
37003701
return Adapter
37013702
.call_nocheck<UrApiKind::urBindlessImagesSignalExternalSemaphoreExp>(
37023703
MQueue->getHandleRef(), SemSignal->getExternalSemaphore(),
3703-
OptSignalValue.has_value(), SignalValue, 0, nullptr, nullptr);
3704+
OptSignalValue.has_value(), SignalValue, RawEvents.size(),
3705+
RawEvents.data(), Event);
37043706
}
37053707
case CGType::AsyncAlloc: {
37063708
// NO-OP. Async alloc calls adapter immediately in order to return a valid

‎sycl/unittests/Extensions/BindlessImages/CMakeLists.txt

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,3 @@
1+
add_sycl_unittest(BindlessImagesExtensionTests OBJECT
2+
Semaphores.cpp
3+
)

‎sycl/unittests/Extensions/BindlessImages/Semaphores.cpp

Lines changed: 161 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,161 @@
1+
#include <helpers/UrMock.hpp>
2+
3+
#include <gtest/gtest.h>
4+
5+
#include <detail/event_impl.hpp>
6+
#include <sycl/detail/core.hpp>
7+
#include <sycl/ext/oneapi/bindless_images.hpp>
8+
#include <sycl/ext/oneapi/bindless_images_interop.hpp>
9+
#include <sycl/queue.hpp>
10+
11+
namespace syclexp = sycl::ext::oneapi::experimental;
12+
13+
constexpr uint64_t WaitValue = 42;
14+
constexpr uint64_t SignalValue = 24;
15+
16+
thread_local int urBindlessImagesWaitExternalSemaphoreExp_counter = 0;
17+
thread_local bool urBindlessImagesWaitExternalSemaphoreExp_expectHasWaitValue =
18+
false;
19+
inline ur_result_t
20+
urBindlessImagesWaitExternalSemaphoreExp_replace(void *pParams) {
21+
++urBindlessImagesWaitExternalSemaphoreExp_counter;
22+
ur_bindless_images_wait_external_semaphore_exp_params_t Params =
23+
*reinterpret_cast<
24+
ur_bindless_images_wait_external_semaphore_exp_params_t *>(pParams);
25+
EXPECT_EQ(*Params.phasWaitValue,
26+
urBindlessImagesWaitExternalSemaphoreExp_expectHasWaitValue);
27+
if (urBindlessImagesWaitExternalSemaphoreExp_expectHasWaitValue) {
28+
EXPECT_EQ(*Params.pwaitValue, WaitValue);
29+
}
30+
EXPECT_EQ(*Params.pnumEventsInWaitList, uint32_t{0});
31+
EXPECT_EQ(*Params.pphEventWaitList, nullptr);
32+
EXPECT_NE(*Params.pphEvent, nullptr);
33+
return UR_RESULT_SUCCESS;
34+
}
35+
36+
thread_local int urBindlessImagesSignalExternalSemaphoreExp_counter = 0;
37+
thread_local bool
38+
urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = false;
39+
thread_local uint32_t
40+
urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 0;
41+
inline ur_result_t
42+
urBindlessImagesSignalExternalSemaphoreExp_replace(void *pParams) {
43+
++urBindlessImagesSignalExternalSemaphoreExp_counter;
44+
ur_bindless_images_signal_external_semaphore_exp_params_t Params =
45+
*reinterpret_cast<
46+
ur_bindless_images_signal_external_semaphore_exp_params_t *>(pParams);
47+
EXPECT_EQ(*Params.phasSignalValue,
48+
urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue);
49+
if (urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue) {
50+
EXPECT_EQ(*Params.psignalValue, SignalValue);
51+
}
52+
EXPECT_EQ(*Params.pnumEventsInWaitList,
53+
urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents);
54+
if (urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents) {
55+
EXPECT_NE(*Params.pphEventWaitList, nullptr);
56+
} else {
57+
EXPECT_EQ(*Params.pphEventWaitList, nullptr);
58+
}
59+
EXPECT_NE(*Params.pphEvent, nullptr);
60+
return UR_RESULT_SUCCESS;
61+
}
62+
63+
TEST(BindlessImagesExtensionTests, ExternalSemaphoreWait) {
64+
sycl::unittest::UrMock<> Mock;
65+
mock::getCallbacks().set_replace_callback(
66+
"urBindlessImagesWaitExternalSemaphoreExp",
67+
&urBindlessImagesWaitExternalSemaphoreExp_replace);
68+
urBindlessImagesWaitExternalSemaphoreExp_counter = 0;
69+
70+
sycl::queue Q;
71+
72+
// Create a dummy external semaphore and set the raw handle to some dummy.
73+
// The mock implementation should never access the handle, so this is safe.
74+
int DummyInt = 0;
75+
syclexp::external_semaphore DummySemaphore{};
76+
DummySemaphore.raw_handle =
77+
reinterpret_cast<ur_exp_external_semaphore_handle_t>(&DummyInt);
78+
79+
DummySemaphore.handle_type =
80+
syclexp::external_semaphore_handle_type::opaque_fd;
81+
82+
urBindlessImagesWaitExternalSemaphoreExp_expectHasWaitValue = false;
83+
Q.ext_oneapi_wait_external_semaphore(DummySemaphore);
84+
EXPECT_EQ(urBindlessImagesWaitExternalSemaphoreExp_counter, 1);
85+
86+
DummySemaphore.handle_type =
87+
syclexp::external_semaphore_handle_type::timeline_fd;
88+
89+
urBindlessImagesWaitExternalSemaphoreExp_expectHasWaitValue = true;
90+
Q.ext_oneapi_wait_external_semaphore(DummySemaphore, WaitValue);
91+
EXPECT_EQ(urBindlessImagesWaitExternalSemaphoreExp_counter, 2);
92+
}
93+
94+
TEST(BindlessImagesExtensionTests, ExternalSemaphoreSignal) {
95+
sycl::unittest::UrMock<> Mock;
96+
mock::getCallbacks().set_replace_callback(
97+
"urBindlessImagesSignalExternalSemaphoreExp",
98+
&urBindlessImagesSignalExternalSemaphoreExp_replace);
99+
urBindlessImagesSignalExternalSemaphoreExp_counter = 0;
100+
101+
sycl::queue Q;
102+
103+
// Create a dummy external semaphore and set the raw handle to some dummy.
104+
// The mock implementation should never access the handle, so this is safe.
105+
int DummyInt1 = 0, DummyInt2 = 0;
106+
syclexp::external_semaphore DummySemaphore{};
107+
DummySemaphore.raw_handle =
108+
reinterpret_cast<ur_exp_external_semaphore_handle_t>(&DummyInt1);
109+
110+
// We create dummy events with dummy UR handles to make the runtime think we
111+
// pass actual device events.
112+
auto DummyEventImpl1 = sycl::detail::event_impl::create_device_event(
113+
*sycl::detail::getSyclObjImpl(Q));
114+
auto DummyEventImpl2 = sycl::detail::event_impl::create_device_event(
115+
*sycl::detail::getSyclObjImpl(Q));
116+
DummyEventImpl1->setHandle(reinterpret_cast<ur_event_handle_t>(&DummyInt1));
117+
DummyEventImpl2->setHandle(reinterpret_cast<ur_event_handle_t>(&DummyInt2));
118+
sycl::event DummyEvent1 =
119+
sycl::detail::createSyclObjFromImpl<sycl::event>(DummyEventImpl1);
120+
sycl::event DummyEvent2 =
121+
sycl::detail::createSyclObjFromImpl<sycl::event>(DummyEventImpl2);
122+
std::vector<sycl::event> DummyEventList{DummyEvent1, DummyEvent2};
123+
124+
DummySemaphore.handle_type =
125+
syclexp::external_semaphore_handle_type::opaque_fd;
126+
127+
urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = false;
128+
urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 0;
129+
Q.ext_oneapi_signal_external_semaphore(DummySemaphore);
130+
EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 1);
131+
132+
urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = false;
133+
urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 1;
134+
Q.ext_oneapi_signal_external_semaphore(DummySemaphore, DummyEvent1);
135+
EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 2);
136+
137+
urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = false;
138+
urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 2;
139+
Q.ext_oneapi_signal_external_semaphore(DummySemaphore, DummyEventList);
140+
EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 3);
141+
142+
DummySemaphore.handle_type =
143+
syclexp::external_semaphore_handle_type::timeline_fd;
144+
145+
urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = true;
146+
urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 0;
147+
Q.ext_oneapi_signal_external_semaphore(DummySemaphore, SignalValue);
148+
EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 4);
149+
150+
urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = true;
151+
urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 1;
152+
Q.ext_oneapi_signal_external_semaphore(DummySemaphore, SignalValue,
153+
DummyEvent1);
154+
EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 5);
155+
156+
urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = true;
157+
urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 2;
158+
Q.ext_oneapi_signal_external_semaphore(DummySemaphore, SignalValue,
159+
DummyEventList);
160+
EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 6);
161+
}

‎sycl/unittests/Extensions/CMakeLists.txt

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -26,6 +26,7 @@ add_sycl_unittest(ExtensionsTests OBJECT
2626
USMPrefetch.cpp
2727
)
2828

29+
add_subdirectory(BindlessImages)
2930
add_subdirectory(CommandGraph)
3031
add_subdirectory(VirtualFunctions)
3132
add_subdirectory(VirtualMemory)

0 commit comments

Comments
 (0)