Skip to content

Commit 268c89e

Browse files
steffenlarsensteffenlarsen
authored
[SYCL] Cherry-pick of external semaphore fixes (intel#20196)
Cherry picks of bindless images semaphore fixes: - [SYCL][BindlessImages] Fix external semaphore dependencies and return events (intel#20040) (b578d54) - [SYCL][BindlessImages] Fix storing result events for semaphores (intel#20080) (68f3fdf) --------- Signed-off-by: Larsen, Steffen <[email protected]>
1 parent 0d3918e commit 268c89e

File tree

4 files changed

+215
-6
lines changed

4 files changed

+215
-6
lines changed

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

Lines changed: 20 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -3710,10 +3710,17 @@ ur_result_t ExecCGCommand::enqueueImpQueue() {
37103710
auto OptWaitValue = SemWait->getWaitValue();
37113711
uint64_t WaitValue = OptWaitValue.has_value() ? OptWaitValue.value() : 0;
37123712

3713-
return Adapter
3714-
.call_nocheck<UrApiKind::urBindlessImagesWaitExternalSemaphoreExp>(
3713+
if (auto Result = Adapter.call_nocheck<
3714+
UrApiKind::urBindlessImagesWaitExternalSemaphoreExp>(
37153715
MQueue->getHandleRef(), SemWait->getExternalSemaphore(),
3716-
OptWaitValue.has_value(), WaitValue, 0, nullptr, nullptr);
3716+
OptWaitValue.has_value(), WaitValue, RawEvents.size(),
3717+
RawEvents.data(), Event);
3718+
Result != UR_RESULT_SUCCESS)
3719+
return Result;
3720+
3721+
SetEventHandleOrDiscard();
3722+
3723+
return UR_RESULT_SUCCESS;
37173724
}
37183725
case CGType::SemaphoreSignal: {
37193726
assert(MQueue &&
@@ -3723,10 +3730,17 @@ ur_result_t ExecCGCommand::enqueueImpQueue() {
37233730
auto OptSignalValue = SemSignal->getSignalValue();
37243731
uint64_t SignalValue =
37253732
OptSignalValue.has_value() ? OptSignalValue.value() : 0;
3726-
return Adapter
3727-
.call_nocheck<UrApiKind::urBindlessImagesSignalExternalSemaphoreExp>(
3733+
if (auto Result = Adapter.call_nocheck<
3734+
UrApiKind::urBindlessImagesSignalExternalSemaphoreExp>(
37283735
MQueue->getHandleRef(), SemSignal->getExternalSemaphore(),
3729-
OptSignalValue.has_value(), SignalValue, 0, nullptr, nullptr);
3736+
OptSignalValue.has_value(), SignalValue, RawEvents.size(),
3737+
RawEvents.data(), Event);
3738+
Result != UR_RESULT_SUCCESS)
3739+
return Result;
3740+
3741+
SetEventHandleOrDiscard();
3742+
3743+
return UR_RESULT_SUCCESS;
37303744
}
37313745
case CGType::AsyncAlloc: {
37323746
// 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: 191 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,191 @@
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+
thread_local ur_event_handle_t
20+
urBindlessImagesWaitExternalSemaphoreExp_lastEvent = nullptr;
21+
inline ur_result_t
22+
urBindlessImagesWaitExternalSemaphoreExp_replace(void *pParams) {
23+
++urBindlessImagesWaitExternalSemaphoreExp_counter;
24+
ur_bindless_images_wait_external_semaphore_exp_params_t Params =
25+
*reinterpret_cast<
26+
ur_bindless_images_wait_external_semaphore_exp_params_t *>(pParams);
27+
EXPECT_EQ(*Params.phasWaitValue,
28+
urBindlessImagesWaitExternalSemaphoreExp_expectHasWaitValue);
29+
if (urBindlessImagesWaitExternalSemaphoreExp_expectHasWaitValue) {
30+
EXPECT_EQ(*Params.pwaitValue, WaitValue);
31+
}
32+
EXPECT_EQ(*Params.pnumEventsInWaitList, uint32_t{0});
33+
EXPECT_EQ(*Params.pphEventWaitList, nullptr);
34+
EXPECT_NE(*Params.pphEvent, nullptr);
35+
if (*Params.pphEvent) {
36+
urBindlessImagesWaitExternalSemaphoreExp_lastEvent =
37+
mock::createDummyHandle<ur_event_handle_t>();
38+
**Params.pphEvent = urBindlessImagesWaitExternalSemaphoreExp_lastEvent;
39+
}
40+
return UR_RESULT_SUCCESS;
41+
}
42+
43+
thread_local int urBindlessImagesSignalExternalSemaphoreExp_counter = 0;
44+
thread_local bool
45+
urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = false;
46+
thread_local uint32_t
47+
urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 0;
48+
thread_local ur_event_handle_t
49+
urBindlessImagesSignalExternalSemaphoreExp_lastEvent = nullptr;
50+
inline ur_result_t
51+
urBindlessImagesSignalExternalSemaphoreExp_replace(void *pParams) {
52+
++urBindlessImagesSignalExternalSemaphoreExp_counter;
53+
ur_bindless_images_signal_external_semaphore_exp_params_t Params =
54+
*reinterpret_cast<
55+
ur_bindless_images_signal_external_semaphore_exp_params_t *>(pParams);
56+
EXPECT_EQ(*Params.phasSignalValue,
57+
urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue);
58+
if (urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue) {
59+
EXPECT_EQ(*Params.psignalValue, SignalValue);
60+
}
61+
EXPECT_EQ(*Params.pnumEventsInWaitList,
62+
urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents);
63+
if (urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents) {
64+
EXPECT_NE(*Params.pphEventWaitList, nullptr);
65+
} else {
66+
EXPECT_EQ(*Params.pphEventWaitList, nullptr);
67+
}
68+
EXPECT_NE(*Params.pphEvent, nullptr);
69+
if (*Params.pphEvent) {
70+
urBindlessImagesSignalExternalSemaphoreExp_lastEvent =
71+
mock::createDummyHandle<ur_event_handle_t>();
72+
**Params.pphEvent = urBindlessImagesSignalExternalSemaphoreExp_lastEvent;
73+
}
74+
return UR_RESULT_SUCCESS;
75+
}
76+
77+
TEST(BindlessImagesExtensionTests, ExternalSemaphoreWait) {
78+
sycl::unittest::UrMock<> Mock;
79+
mock::getCallbacks().set_replace_callback(
80+
"urBindlessImagesWaitExternalSemaphoreExp",
81+
&urBindlessImagesWaitExternalSemaphoreExp_replace);
82+
urBindlessImagesWaitExternalSemaphoreExp_counter = 0;
83+
84+
sycl::queue Q;
85+
86+
// Create a dummy external semaphore and set the raw handle to some dummy.
87+
// The mock implementation should never access the handle, so this is safe.
88+
int DummyInt = 0;
89+
syclexp::external_semaphore DummySemaphore{};
90+
DummySemaphore.raw_handle =
91+
reinterpret_cast<ur_exp_external_semaphore_handle_t>(&DummyInt);
92+
93+
DummySemaphore.handle_type =
94+
syclexp::external_semaphore_handle_type::opaque_fd;
95+
96+
urBindlessImagesWaitExternalSemaphoreExp_expectHasWaitValue = false;
97+
sycl::event E = Q.ext_oneapi_wait_external_semaphore(DummySemaphore);
98+
EXPECT_EQ(urBindlessImagesWaitExternalSemaphoreExp_counter, 1);
99+
EXPECT_EQ(sycl::detail::getSyclObjImpl(E)->getHandle(),
100+
urBindlessImagesWaitExternalSemaphoreExp_lastEvent);
101+
102+
DummySemaphore.handle_type =
103+
syclexp::external_semaphore_handle_type::timeline_fd;
104+
105+
urBindlessImagesWaitExternalSemaphoreExp_expectHasWaitValue = true;
106+
E = Q.ext_oneapi_wait_external_semaphore(DummySemaphore, WaitValue);
107+
EXPECT_EQ(urBindlessImagesWaitExternalSemaphoreExp_counter, 2);
108+
EXPECT_EQ(sycl::detail::getSyclObjImpl(E)->getHandle(),
109+
urBindlessImagesWaitExternalSemaphoreExp_lastEvent);
110+
}
111+
112+
TEST(BindlessImagesExtensionTests, ExternalSemaphoreSignal) {
113+
sycl::unittest::UrMock<> Mock;
114+
mock::getCallbacks().set_replace_callback(
115+
"urBindlessImagesSignalExternalSemaphoreExp",
116+
&urBindlessImagesSignalExternalSemaphoreExp_replace);
117+
urBindlessImagesSignalExternalSemaphoreExp_counter = 0;
118+
119+
sycl::queue Q;
120+
121+
// Create a dummy external semaphore and set the raw handle to some dummy.
122+
// The mock implementation should never access the handle, so this is safe.
123+
int DummyInt1 = 0, DummyInt2 = 0;
124+
syclexp::external_semaphore DummySemaphore{};
125+
DummySemaphore.raw_handle =
126+
reinterpret_cast<ur_exp_external_semaphore_handle_t>(&DummyInt1);
127+
128+
// We create dummy events with dummy UR handles to make the runtime think we
129+
// pass actual device events.
130+
auto DummyEventImpl1 = sycl::detail::event_impl::create_device_event(
131+
*sycl::detail::getSyclObjImpl(Q));
132+
auto DummyEventImpl2 = sycl::detail::event_impl::create_device_event(
133+
*sycl::detail::getSyclObjImpl(Q));
134+
DummyEventImpl1->setHandle(reinterpret_cast<ur_event_handle_t>(&DummyInt1));
135+
DummyEventImpl2->setHandle(reinterpret_cast<ur_event_handle_t>(&DummyInt2));
136+
sycl::event DummyEvent1 =
137+
sycl::detail::createSyclObjFromImpl<sycl::event>(DummyEventImpl1);
138+
sycl::event DummyEvent2 =
139+
sycl::detail::createSyclObjFromImpl<sycl::event>(DummyEventImpl2);
140+
std::vector<sycl::event> DummyEventList{DummyEvent1, DummyEvent2};
141+
142+
DummySemaphore.handle_type =
143+
syclexp::external_semaphore_handle_type::opaque_fd;
144+
145+
urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = false;
146+
urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 0;
147+
sycl::event E = Q.ext_oneapi_signal_external_semaphore(DummySemaphore);
148+
EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 1);
149+
EXPECT_EQ(sycl::detail::getSyclObjImpl(E)->getHandle(),
150+
urBindlessImagesSignalExternalSemaphoreExp_lastEvent);
151+
152+
urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = false;
153+
urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 1;
154+
E = Q.ext_oneapi_signal_external_semaphore(DummySemaphore, DummyEvent1);
155+
EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 2);
156+
EXPECT_EQ(sycl::detail::getSyclObjImpl(E)->getHandle(),
157+
urBindlessImagesSignalExternalSemaphoreExp_lastEvent);
158+
159+
urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = false;
160+
urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 2;
161+
E = Q.ext_oneapi_signal_external_semaphore(DummySemaphore, DummyEventList);
162+
EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 3);
163+
EXPECT_EQ(sycl::detail::getSyclObjImpl(E)->getHandle(),
164+
urBindlessImagesSignalExternalSemaphoreExp_lastEvent);
165+
166+
DummySemaphore.handle_type =
167+
syclexp::external_semaphore_handle_type::timeline_fd;
168+
169+
urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = true;
170+
urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 0;
171+
E = Q.ext_oneapi_signal_external_semaphore(DummySemaphore, SignalValue);
172+
EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 4);
173+
EXPECT_EQ(sycl::detail::getSyclObjImpl(E)->getHandle(),
174+
urBindlessImagesSignalExternalSemaphoreExp_lastEvent);
175+
176+
urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = true;
177+
urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 1;
178+
E = Q.ext_oneapi_signal_external_semaphore(DummySemaphore, SignalValue,
179+
DummyEvent1);
180+
EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 5);
181+
EXPECT_EQ(sycl::detail::getSyclObjImpl(E)->getHandle(),
182+
urBindlessImagesSignalExternalSemaphoreExp_lastEvent);
183+
184+
urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = true;
185+
urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 2;
186+
E = Q.ext_oneapi_signal_external_semaphore(DummySemaphore, SignalValue,
187+
DummyEventList);
188+
EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 6);
189+
EXPECT_EQ(sycl::detail::getSyclObjImpl(E)->getHandle(),
190+
urBindlessImagesSignalExternalSemaphoreExp_lastEvent);
191+
}

‎sycl/unittests/Extensions/CMakeLists.txt‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -24,6 +24,7 @@ add_sycl_unittest(ExtensionsTests OBJECT
2424
RootGroup.cpp
2525
)
2626

27+
add_subdirectory(BindlessImages)
2728
add_subdirectory(CommandGraph)
2829
add_subdirectory(VirtualFunctions)
2930
add_subdirectory(VirtualMemory)

0 commit comments

Comments
 (0)