Skip to content

Commit f4bfb3b

Browse files
steffenlarsensteffenlarsen
committed
[SYCL][BindlessImages] Fix external semaphore dependencies and return events
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 cb03a1b commit f4bfb3b

File tree

4 files changed

+135
-2
lines changed

4 files changed

+135
-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: 127 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,127 @@
1+
#include <helpers/UrMock.hpp>
2+
3+
#include <gtest/gtest.h>
4+
5+
#include <sycl/detail/core.hpp>
6+
#include <sycl/ext/oneapi/bindless_images.hpp>
7+
#include <sycl/ext/oneapi/bindless_images_interop.hpp>
8+
#include <sycl/queue.hpp>
9+
10+
namespace syclexp = sycl::ext::oneapi::experimental;
11+
12+
constexpr uint64_t WaitValue = 42;
13+
constexpr uint64_t SignalValue = 24;
14+
15+
thread_local int urBindlessImagesWaitExternalSemaphoreExp_counter = 0;
16+
thread_local bool urBindlessImagesWaitExternalSemaphoreExp_expectHasWaitValue =
17+
false;
18+
inline ur_result_t
19+
urBindlessImagesWaitExternalSemaphoreExp_replace(void *pParams) {
20+
++urBindlessImagesWaitExternalSemaphoreExp_counter;
21+
ur_bindless_images_wait_external_semaphore_exp_params_t Params =
22+
*reinterpret_cast<
23+
ur_bindless_images_wait_external_semaphore_exp_params_t *>(pParams);
24+
EXPECT_EQ(*Params.phasWaitValue,
25+
urBindlessImagesWaitExternalSemaphoreExp_expectHasWaitValue);
26+
if (urBindlessImagesWaitExternalSemaphoreExp_expectHasWaitValue) {
27+
EXPECT_EQ(*Params.pwaitValue, WaitValue);
28+
}
29+
EXPECT_EQ(*Params.pphEvent, nullptr);
30+
EXPECT_EQ(*Params.pnumEventsInWaitList, uint32_t{0});
31+
EXPECT_NE(*Params.pphEventWaitList, nullptr);
32+
return UR_RESULT_SUCCESS;
33+
}
34+
35+
thread_local int urBindlessImagesSignalExternalSemaphoreExp_counter = 0;
36+
thread_local bool
37+
urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = false;
38+
thread_local uint32_t
39+
urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 0;
40+
inline ur_result_t
41+
urBindlessImagesSignalExternalSemaphoreExp_replace(void *pParams) {
42+
++urBindlessImagesSignalExternalSemaphoreExp_counter;
43+
ur_bindless_images_signal_external_semaphore_exp_params_t Params =
44+
*reinterpret_cast<
45+
ur_bindless_images_signal_external_semaphore_exp_params_t *>(pParams);
46+
EXPECT_EQ(*Params.pphEvent, nullptr);
47+
EXPECT_EQ(*Params.phasSignalValue,
48+
urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue);
49+
if (urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue) {
50+
EXPECT_EQ(*Params.psignalValue, SignalValue);
51+
}
52+
if (urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents) {
53+
EXPECT_NE(*Params.pphEvent, nullptr);
54+
}
55+
56+
else {
57+
EXPECT_EQ(*Params.pphEvent, nullptr);
58+
}
59+
EXPECT_EQ(*Params.pnumEventsInWaitList,
60+
urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents);
61+
EXPECT_NE(*Params.pphEventWaitList, nullptr);
62+
return UR_RESULT_SUCCESS;
63+
}
64+
65+
TEST(BindlessImagesExtensionTests, ExternalSemaphoreWait) {
66+
sycl::unittest::UrMock<> Mock;
67+
mock::getCallbacks().set_replace_callback(
68+
"urBindlessImagesWaitExternalSemaphoreExp",
69+
&urBindlessImagesWaitExternalSemaphoreExp_replace);
70+
urBindlessImagesWaitExternalSemaphoreExp_counter = 0;
71+
72+
sycl::queue Q;
73+
syclexp::external_semaphore DummySemaphore{};
74+
75+
urBindlessImagesWaitExternalSemaphoreExp_expectHasWaitValue = false;
76+
Q.ext_oneapi_wait_external_semaphore(DummySemaphore);
77+
EXPECT_EQ(urBindlessImagesWaitExternalSemaphoreExp_counter, 1);
78+
79+
urBindlessImagesWaitExternalSemaphoreExp_expectHasWaitValue = true;
80+
Q.ext_oneapi_wait_external_semaphore(DummySemaphore, WaitValue);
81+
EXPECT_EQ(urBindlessImagesWaitExternalSemaphoreExp_counter, 2);
82+
}
83+
84+
TEST(BindlessImagesExtensionTests, ExternalSemaphoreSignal) {
85+
sycl::unittest::UrMock<> Mock;
86+
mock::getCallbacks().set_replace_callback(
87+
"urBindlessImagesSignalExternalSemaphoreExp",
88+
&urBindlessImagesSignalExternalSemaphoreExp_replace);
89+
urBindlessImagesSignalExternalSemaphoreExp_counter = 0;
90+
91+
sycl::queue Q;
92+
syclexp::external_semaphore DummySemaphore{};
93+
94+
urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = false;
95+
urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 0;
96+
Q.ext_oneapi_signal_external_semaphore(DummySemaphore);
97+
EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 1);
98+
99+
urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = true;
100+
urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 0;
101+
Q.ext_oneapi_signal_external_semaphore(DummySemaphore, SignalValue);
102+
EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 2);
103+
104+
urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = false;
105+
urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 1;
106+
Q.ext_oneapi_signal_external_semaphore(DummySemaphore, sycl::event{});
107+
EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 3);
108+
109+
urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = true;
110+
urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 1;
111+
Q.ext_oneapi_signal_external_semaphore(DummySemaphore, SignalValue,
112+
sycl::event{});
113+
EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 4);
114+
115+
std::vector<sycl::event> DummyEventList(2);
116+
117+
urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = false;
118+
urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 2;
119+
Q.ext_oneapi_signal_external_semaphore(DummySemaphore, DummyEventList);
120+
EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 5);
121+
122+
urBindlessImagesSignalExternalSemaphoreExp_expectHasSignalValue = true;
123+
urBindlessImagesSignalExternalSemaphoreExp_expectedNumWaitEvents = 2;
124+
Q.ext_oneapi_signal_external_semaphore(DummySemaphore, SignalValue,
125+
DummyEventList);
126+
EXPECT_EQ(urBindlessImagesSignalExternalSemaphoreExp_counter, 6);
127+
}

‎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)