Skip to content

Commit 6f5c5ae

Browse files
authored
Merge pull request #1586 from RossBrunton/ross/eventtest
[CTS] Added test for sequencing events
2 parents 4f8aa0d + 759840e commit 6f5c5ae

File tree

5 files changed

+202
-2
lines changed

5 files changed

+202
-2
lines changed

test/conformance/device_code/CMakeLists.txt

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -158,6 +158,7 @@ add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/saxpy_usm.cpp)
158158
add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/indexers_usm.cpp)
159159
add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/build_failure.cpp)
160160
add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/fixed_wg_size.cpp)
161+
add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/sequence.cpp)
161162
add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/standard_types.cpp)
162163
add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/subgroup.cpp)
163164
add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/linker_error.cpp)
Lines changed: 39 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,39 @@
1+
// Copyright (C) 2024 Intel Corporation
2+
// Part of the Unified-Runtime Project, under the Apache License v2.0 with LLVM Exceptions.
3+
// See LICENSE.TXT
4+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
5+
6+
#include <stdint.h>
7+
#include <sycl/sycl.hpp>
8+
9+
class Add;
10+
class Mul;
11+
12+
int main() {
13+
sycl::queue deviceQueue;
14+
uint32_t val = 0;
15+
16+
auto buff = sycl::buffer<uint32_t>(&val, 1);
17+
18+
deviceQueue.submit([&](sycl::handler &cgh) {
19+
auto acc = buff.get_access<sycl::access::mode::read_write>(cgh);
20+
cgh.single_task<Add>([=]() {
21+
for (uint32_t i = 0; i < 1000; i++) {
22+
volatile uint32_t tmp = acc[0];
23+
acc[0] = tmp + 1;
24+
}
25+
});
26+
});
27+
28+
deviceQueue.submit([&](sycl::handler &cgh) {
29+
auto acc = buff.get_access<sycl::access::mode::read_write>(cgh);
30+
cgh.single_task<Mul>([=]() {
31+
for (uint32_t i = 0; i < 2; i++) {
32+
volatile uint32_t tmp = acc[0];
33+
acc[0] = tmp * 2;
34+
}
35+
});
36+
});
37+
38+
return 0;
39+
}

test/conformance/enqueue/enqueue_adapter_native_cpu.match

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -18,6 +18,10 @@
1818
{{OPT}}urEnqueueEventsWaitTest.InvalidNullPtrEventWaitList/SYCL_NATIVE_CPU___SYCL_Native_CPU__{{.*}}
1919
{{OPT}}urEnqueueEventsWaitWithBarrierTest.Success/SYCL_NATIVE_CPU___SYCL_Native_CPU__{{.*}}
2020
{{OPT}}urEnqueueEventsWaitWithBarrierTest.InvalidNullPtrEventWaitList/SYCL_NATIVE_CPU___SYCL_Native_CPU__{{.*}}
21+
urEnqueueEventsWaitWithBarrierOrderingTest.SuccessEventDependenciesBarrierOnly/SYCL_NATIVE_CPU___SYCL_Native_CPU__{{.*}}_
22+
urEnqueueEventsWaitWithBarrierOrderingTest.SuccessEventDependenciesLaunchOnly/SYCL_NATIVE_CPU___SYCL_Native_CPU__{{.*}}_
23+
urEnqueueEventsWaitWithBarrierOrderingTest.SuccessEventDependencies/SYCL_NATIVE_CPU___SYCL_Native_CPU__{{.*}}_
24+
urEnqueueEventsWaitWithBarrierOrderingTest.SuccessNonEventDependencies/SYCL_NATIVE_CPU___SYCL_Native_CPU__{{.*}}_
2125
{{OPT}}urEnqueueKernelLaunchTest.Success/SYCL_NATIVE_CPU___SYCL_Native_CPU__{{.*}}
2226
{{OPT}}urEnqueueKernelLaunchTest.InvalidNullHandleQueue/SYCL_NATIVE_CPU___SYCL_Native_CPU__{{.*}}
2327
{{OPT}}urEnqueueKernelLaunchTest.InvalidNullHandleKernel/SYCL_NATIVE_CPU___SYCL_Native_CPU__{{.*}}

test/conformance/enqueue/urEnqueueEventsWaitWithBarrier.cpp

Lines changed: 153 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// Copyright (C) 2023 Intel Corporation
1+
// Copyright (C) 2024 Intel Corporation
22
// Part of the Unified-Runtime Project, under the Apache License v2.0 with LLVM Exceptions.
33
// See LICENSE.TXT
44
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
@@ -36,6 +36,32 @@ struct urEnqueueEventsWaitWithBarrierTest : uur::urMultiQueueTest {
3636

3737
UUR_INSTANTIATE_DEVICE_TEST_SUITE_P(urEnqueueEventsWaitWithBarrierTest);
3838

39+
struct urEnqueueEventsWaitWithBarrierOrderingTest : uur::urProgramTest {
40+
void SetUp() override {
41+
program_name = "sequence";
42+
UUR_RETURN_ON_FATAL_FAILURE(urProgramTest::SetUp());
43+
ASSERT_SUCCESS(urProgramBuild(context, program, nullptr));
44+
45+
ASSERT_SUCCESS(urMemBufferCreate(context, UR_MEM_FLAG_READ_WRITE,
46+
sizeof(uint32_t), nullptr, &buffer));
47+
48+
auto entry_points =
49+
uur::KernelsEnvironment::instance->GetEntryPointNames(program_name);
50+
std::cout << entry_points[0];
51+
52+
ASSERT_SUCCESS(urKernelCreate(program, "_ZTS3Add", &add_kernel));
53+
ASSERT_SUCCESS(urKernelCreate(program, "_ZTS3Mul", &mul_kernel));
54+
}
55+
56+
void TearDown() override { uur::urProgramTest::TearDown(); }
57+
58+
ur_kernel_handle_t add_kernel;
59+
ur_kernel_handle_t mul_kernel;
60+
ur_mem_handle_t buffer = nullptr;
61+
};
62+
63+
UUR_INSTANTIATE_DEVICE_TEST_SUITE_P(urEnqueueEventsWaitWithBarrierOrderingTest);
64+
3965
TEST_P(urEnqueueEventsWaitWithBarrierTest, Success) {
4066
ur_event_handle_t event1 = nullptr;
4167
ur_event_handle_t waitEvent = nullptr;
@@ -97,3 +123,129 @@ TEST_P(urEnqueueEventsWaitWithBarrierTest, InvalidNullPtrEventWaitList) {
97123

98124
ASSERT_SUCCESS(urEventRelease(validEvent));
99125
}
126+
127+
TEST_P(urEnqueueEventsWaitWithBarrierOrderingTest,
128+
SuccessEventDependenciesBarrierOnly) {
129+
constexpr size_t offset = 0;
130+
constexpr size_t count = 1;
131+
ur_event_handle_t event;
132+
133+
uur::KernelLaunchHelper addHelper(platform, context, add_kernel, queue);
134+
uur::KernelLaunchHelper mulHelper(platform, context, mul_kernel, queue);
135+
136+
addHelper.SetBuffer1DArg(buffer, nullptr);
137+
mulHelper.SetBuffer1DArg(buffer, nullptr);
138+
139+
for (size_t i = 0; i < 10; i++) {
140+
constexpr uint32_t ONE = 1;
141+
urEnqueueMemBufferWrite(queue, buffer, true, 0, sizeof(uint32_t), &ONE,
142+
0, nullptr, &event);
143+
EXPECT_SUCCESS(
144+
urEnqueueEventsWaitWithBarrier(queue, 1, &event, nullptr));
145+
EXPECT_SUCCESS(urEnqueueKernelLaunch(queue, add_kernel, 1, &offset,
146+
&count, nullptr, 0, nullptr,
147+
&event));
148+
EXPECT_SUCCESS(
149+
urEnqueueEventsWaitWithBarrier(queue, 1, &event, nullptr));
150+
EXPECT_SUCCESS(urEnqueueKernelLaunch(queue, mul_kernel, 1, &offset,
151+
&count, nullptr, 0, nullptr,
152+
&event));
153+
EXPECT_SUCCESS(
154+
urEnqueueEventsWaitWithBarrier(queue, 1, &event, nullptr));
155+
addHelper.ValidateBuffer(buffer, sizeof(uint32_t), 4004);
156+
}
157+
}
158+
159+
TEST_P(urEnqueueEventsWaitWithBarrierOrderingTest,
160+
SuccessEventDependenciesLaunchOnly) {
161+
constexpr size_t offset = 0;
162+
constexpr size_t count = 1;
163+
ur_event_handle_t event;
164+
165+
uur::KernelLaunchHelper addHelper(platform, context, add_kernel, queue);
166+
uur::KernelLaunchHelper mulHelper(platform, context, mul_kernel, queue);
167+
168+
addHelper.SetBuffer1DArg(buffer, nullptr);
169+
mulHelper.SetBuffer1DArg(buffer, nullptr);
170+
171+
for (size_t i = 0; i < 10; i++) {
172+
constexpr uint32_t ONE = 1;
173+
urEnqueueMemBufferWrite(queue, buffer, true, 0, sizeof(uint32_t), &ONE,
174+
0, nullptr, nullptr);
175+
EXPECT_SUCCESS(
176+
urEnqueueEventsWaitWithBarrier(queue, 0, nullptr, &event));
177+
EXPECT_SUCCESS(urEnqueueKernelLaunch(queue, add_kernel, 1, &offset,
178+
&count, nullptr, 1, &event,
179+
nullptr));
180+
EXPECT_SUCCESS(
181+
urEnqueueEventsWaitWithBarrier(queue, 0, nullptr, &event));
182+
EXPECT_SUCCESS(urEnqueueKernelLaunch(queue, mul_kernel, 1, &offset,
183+
&count, nullptr, 1, &event,
184+
nullptr));
185+
EXPECT_SUCCESS(
186+
urEnqueueEventsWaitWithBarrier(queue, 0, nullptr, &event));
187+
addHelper.ValidateBuffer(buffer, sizeof(uint32_t), 4004);
188+
}
189+
}
190+
191+
TEST_P(urEnqueueEventsWaitWithBarrierOrderingTest, SuccessEventDependencies) {
192+
constexpr size_t offset = 0;
193+
constexpr size_t count = 1;
194+
ur_event_handle_t event[6];
195+
196+
uur::KernelLaunchHelper addHelper(platform, context, add_kernel, queue);
197+
uur::KernelLaunchHelper mulHelper(platform, context, mul_kernel, queue);
198+
199+
addHelper.SetBuffer1DArg(buffer, nullptr);
200+
mulHelper.SetBuffer1DArg(buffer, nullptr);
201+
202+
for (size_t i = 0; i < 10; i++) {
203+
constexpr uint32_t ONE = 1;
204+
urEnqueueMemBufferWrite(queue, buffer, true, 0, sizeof(uint32_t), &ONE,
205+
0, nullptr, &event[0]);
206+
EXPECT_SUCCESS(
207+
urEnqueueEventsWaitWithBarrier(queue, 1, &event[0], &event[1]));
208+
EXPECT_SUCCESS(urEnqueueKernelLaunch(queue, add_kernel, 1, &offset,
209+
&count, nullptr, 1, &event[1],
210+
&event[2]));
211+
EXPECT_SUCCESS(
212+
urEnqueueEventsWaitWithBarrier(queue, 1, &event[2], &event[3]));
213+
EXPECT_SUCCESS(urEnqueueKernelLaunch(queue, mul_kernel, 1, &offset,
214+
&count, nullptr, 1, &event[3],
215+
&event[4]));
216+
EXPECT_SUCCESS(
217+
urEnqueueEventsWaitWithBarrier(queue, 1, &event[4], &event[5]));
218+
addHelper.ValidateBuffer(buffer, sizeof(uint32_t), 4004);
219+
}
220+
}
221+
222+
TEST_P(urEnqueueEventsWaitWithBarrierOrderingTest,
223+
SuccessNonEventDependencies) {
224+
constexpr size_t offset = 0;
225+
constexpr size_t count = 1;
226+
227+
uur::KernelLaunchHelper addHelper(platform, context, add_kernel, queue);
228+
uur::KernelLaunchHelper mulHelper(platform, context, mul_kernel, queue);
229+
230+
addHelper.SetBuffer1DArg(buffer, nullptr);
231+
mulHelper.SetBuffer1DArg(buffer, nullptr);
232+
233+
for (size_t i = 0; i < 10; i++) {
234+
constexpr uint32_t ONE = 1;
235+
urEnqueueMemBufferWrite(queue, buffer, true, 0, sizeof(uint32_t), &ONE,
236+
0, nullptr, nullptr);
237+
EXPECT_SUCCESS(
238+
urEnqueueEventsWaitWithBarrier(queue, 0, nullptr, nullptr));
239+
EXPECT_SUCCESS(urEnqueueKernelLaunch(queue, add_kernel, 1, &offset,
240+
&count, nullptr, 0, nullptr,
241+
nullptr));
242+
EXPECT_SUCCESS(
243+
urEnqueueEventsWaitWithBarrier(queue, 0, nullptr, nullptr));
244+
EXPECT_SUCCESS(urEnqueueKernelLaunch(queue, mul_kernel, 1, &offset,
245+
&count, nullptr, 0, nullptr,
246+
nullptr));
247+
EXPECT_SUCCESS(
248+
urEnqueueEventsWaitWithBarrier(queue, 0, nullptr, nullptr));
249+
addHelper.ValidateBuffer(buffer, sizeof(uint32_t), 4004);
250+
}
251+
}

test/conformance/testing/include/uur/fixtures.h

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1314,6 +1314,11 @@ struct KernelLaunchHelper {
13141314
sizeof(zero), 0, size, 0, nullptr,
13151315
nullptr));
13161316
ASSERT_SUCCESS(urQueueFinish(queue));
1317+
SetBuffer1DArg(mem_handle, buffer_index);
1318+
*out_buffer = mem_handle;
1319+
}
1320+
1321+
void SetBuffer1DArg(ur_mem_handle_t mem_handle, size_t *buffer_index) {
13171322
ASSERT_SUCCESS(urKernelSetArgMemObj(kernel, current_arg_index, nullptr,
13181323
mem_handle));
13191324
if (buffer_index) {
@@ -1350,7 +1355,6 @@ struct KernelLaunchHelper {
13501355
&accessor));
13511356
current_arg_index += 2;
13521357
}
1353-
*out_buffer = mem_handle;
13541358
}
13551359

13561360
template <class T> void AddPodArg(T data) {

0 commit comments

Comments
 (0)