From e108e6c95cfafad4984e255ae78e204834b01817 Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Mon, 22 Sep 2025 23:34:01 +0200 Subject: [PATCH 1/5] [SYCL] Don't return last event in ext_oneapi_submit_barrier --- sycl/source/queue.cpp | 29 ----------------------------- 1 file changed, 29 deletions(-) diff --git a/sycl/source/queue.cpp b/sycl/source/queue.cpp index b73c24091d4ed..e13ff1b377598 100644 --- a/sycl/source/queue.cpp +++ b/sycl/source/queue.cpp @@ -334,24 +334,6 @@ void queue::wait_and_throw_proxy(const detail::code_location &CodeLoc) { impl->wait_and_throw(CodeLoc); } -static event -getBarrierEventForInorderQueueHelper(detail::queue_impl &QueueImpl) { - // This function should not be called when a queue is recording to a graph, - // as a graph can record from multiple queues and we cannot guarantee the - // last node added by an in-order queue will be the last node added to the - // graph. - assert(!QueueImpl.hasCommandGraph() && - "Should not be called in on graph recording."); - - sycl::detail::optional LastEvent = QueueImpl.getLastEvent(); - if (LastEvent) - return *LastEvent; - - // If there was no last event, we create an empty one. - return detail::createSyclObjFromImpl( - detail::event_impl::create_default_event()); -} - /// Prevents any commands submitted afterward to this queue from executing /// until all commands previously submitted to this queue have entered the /// complete state. @@ -374,17 +356,6 @@ event queue::ext_oneapi_submit_barrier(const detail::code_location &CodeLoc) { /// group is being enqueued on. event queue::ext_oneapi_submit_barrier(const std::vector &WaitList, const detail::code_location &CodeLoc) { - bool AllEventsEmptyOrNop = std::all_of( - begin(WaitList), end(WaitList), [&](const event &Event) -> bool { - detail::event_impl &EventImpl = *detail::getSyclObjImpl(Event); - return (EventImpl.isDefaultConstructed() || EventImpl.isNOP()) && - !EventImpl.hasCommandGraph(); - }); - if (is_in_order() && !impl->hasCommandGraph() && !impl->MIsProfilingEnabled && - AllEventsEmptyOrNop) { - return getBarrierEventForInorderQueueHelper(*impl); - } - if (WaitList.empty()) return submit([=](handler &CGH) { CGH.ext_oneapi_barrier(); }, CodeLoc); else From 3e3507ad12244faafaa5e48fbd52f22753c631a7 Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Tue, 23 Sep 2025 00:30:15 +0200 Subject: [PATCH 2/5] Update test --- .../InorderQueue/in_order_ext_oneapi_submit_barrier.cpp | 9 +++------ 1 file changed, 3 insertions(+), 6 deletions(-) diff --git a/sycl/test-e2e/InorderQueue/in_order_ext_oneapi_submit_barrier.cpp b/sycl/test-e2e/InorderQueue/in_order_ext_oneapi_submit_barrier.cpp index f4396e6b31b73..a357c94c36cd2 100644 --- a/sycl/test-e2e/InorderQueue/in_order_ext_oneapi_submit_barrier.cpp +++ b/sycl/test-e2e/InorderQueue/in_order_ext_oneapi_submit_barrier.cpp @@ -53,12 +53,9 @@ int main() { std::cout << "Test 2" << std::endl; *Res = 0; - auto Event1 = Q.submit( - [&](sycl::handler &CGH) { CGH.host_task([&] { *Res += 1; }); }); - auto BarrierEvent1 = Q.ext_oneapi_submit_barrier(); - assert(checkBarrierEvent(Q.get_backend(), Event1, BarrierEvent1, - false /* host tasks used */)); - auto Event2 = Q.submit([&](sycl::handler &CGH) { CGH.fill(Res, 10, 1); }); + Q.submit([&](sycl::handler &CGH) { CGH.host_task([&] { *Res += 1; }); }); + Q.ext_oneapi_submit_barrier(); + Q.submit([&](sycl::handler &CGH) { CGH.fill(Res, 10, 1); }); Q.wait(); assert(*Res == 10); From 9620fff1bc55a68d81cf565c5724743f145a282b Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Tue, 23 Sep 2025 01:18:56 +0200 Subject: [PATCH 3/5] Remove empty queue test --- .../in_order_ext_oneapi_submit_barrier.cpp | 12 ------------ 1 file changed, 12 deletions(-) diff --git a/sycl/test-e2e/InorderQueue/in_order_ext_oneapi_submit_barrier.cpp b/sycl/test-e2e/InorderQueue/in_order_ext_oneapi_submit_barrier.cpp index a357c94c36cd2..f2b1fba5b9359 100644 --- a/sycl/test-e2e/InorderQueue/in_order_ext_oneapi_submit_barrier.cpp +++ b/sycl/test-e2e/InorderQueue/in_order_ext_oneapi_submit_barrier.cpp @@ -60,18 +60,6 @@ int main() { Q.wait(); assert(*Res == 10); } - - { - // Test cast 3 - empty queue. - std::cout << "Test 3" << std::endl; - sycl::queue EmptyQ({sycl::property::queue::in_order{}}); - auto BarrierEvent = EmptyQ.ext_oneapi_submit_barrier(); - assert( - BarrierEvent.get_info() == - sycl::info::event_command_status::complete); - BarrierEvent.wait(); - } - { // Test cast 4 - graph. sycl::queue GQueue{sycl::property::queue::in_order{}}; From eddffb34f25f7a2e3bb391e7a3cfa050eda5b06c Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Wed, 24 Sep 2025 03:09:17 +0200 Subject: [PATCH 4/5] Filter out empty events. --- sycl/source/queue.cpp | 12 +++- .../Regression/ext_oneapi_barrier_opt.cpp | 48 --------------- sycl/unittests/Extensions/CMakeLists.txt | 1 + .../Extensions/ExtOneapiBarrierOpt.cpp | 60 +++++++++++++++++++ 4 files changed, 72 insertions(+), 49 deletions(-) delete mode 100644 sycl/test-e2e/Regression/ext_oneapi_barrier_opt.cpp create mode 100644 sycl/unittests/Extensions/ExtOneapiBarrierOpt.cpp diff --git a/sycl/source/queue.cpp b/sycl/source/queue.cpp index e13ff1b377598..9fc071cc87ba6 100644 --- a/sycl/source/queue.cpp +++ b/sycl/source/queue.cpp @@ -356,7 +356,17 @@ event queue::ext_oneapi_submit_barrier(const detail::code_location &CodeLoc) { /// group is being enqueued on. event queue::ext_oneapi_submit_barrier(const std::vector &WaitList, const detail::code_location &CodeLoc) { - if (WaitList.empty()) + + // If waitlist contains only empty, default constructed events, ignore + // them. + bool AllEventsEmptyOrNop = std::all_of( + begin(WaitList), end(WaitList), [&](const event &Event) -> bool { + detail::event_impl &EventImpl = *detail::getSyclObjImpl(Event); + return (EventImpl.isDefaultConstructed() || EventImpl.isNOP()) && + !EventImpl.hasCommandGraph(); + }); + + if (WaitList.empty() || AllEventsEmptyOrNop) return submit([=](handler &CGH) { CGH.ext_oneapi_barrier(); }, CodeLoc); else return submit([=](handler &CGH) { CGH.ext_oneapi_barrier(WaitList); }, diff --git a/sycl/test-e2e/Regression/ext_oneapi_barrier_opt.cpp b/sycl/test-e2e/Regression/ext_oneapi_barrier_opt.cpp deleted file mode 100644 index 2275c35616945..0000000000000 --- a/sycl/test-e2e/Regression/ext_oneapi_barrier_opt.cpp +++ /dev/null @@ -1,48 +0,0 @@ -// RUN: %{build} %threads_lib -o %t.out -// RUN: %{run} %t.out - -// Check that ext_oneapi_submit_barrier works fine in the scenarios -// when provided waitlist consists of only empty events. - -#include -#include -#include - -#include - -#include - -static constexpr int niter = 1024; -static constexpr int nthreads = 2; - -std::array mutexes; -std::array, nthreads> events; - -void threadFunction(int tid) { - sycl::device dev; - std::cout << dev.get_info() << std::endl; - sycl::context ctx{dev}; - sycl::queue q1{ctx, dev, {sycl::property::queue::in_order()}}; - sycl::queue q2{ctx, dev, {sycl::property::queue::in_order()}}; - for (int i = 0; i < niter; i++) { - sycl::event ev1 = q1.ext_oneapi_submit_barrier(); - q2.ext_oneapi_submit_barrier({ev1}); - sycl::event ev2 = q2.ext_oneapi_submit_barrier(); - q1.ext_oneapi_submit_barrier({ev2}); - } -} - -int main() { - std::array threads; - - for (int i = 0; i < nthreads; i++) { - threads[i] = std::thread{threadFunction, i}; - } - - for (int i = 0; i < nthreads; i++) { - threads[i].join(); - } - std::cout << "All threads have finished." << std::endl; - - return 0; -} diff --git a/sycl/unittests/Extensions/CMakeLists.txt b/sycl/unittests/Extensions/CMakeLists.txt index 32d1b34a99311..311b411985a40 100644 --- a/sycl/unittests/Extensions/CMakeLists.txt +++ b/sycl/unittests/Extensions/CMakeLists.txt @@ -13,6 +13,7 @@ add_sycl_unittest(ExtensionsTests OBJECT OneAPIProd.cpp EnqueueFunctionsEvents.cpp EnqueueFunctionsPrefetch.cpp + ExtOneapiBarrierOpt.cpp ProfilingTag.cpp KernelProperties.cpp NoDeviceIPVersion.cpp diff --git a/sycl/unittests/Extensions/ExtOneapiBarrierOpt.cpp b/sycl/unittests/Extensions/ExtOneapiBarrierOpt.cpp new file mode 100644 index 0000000000000..295273626651f --- /dev/null +++ b/sycl/unittests/Extensions/ExtOneapiBarrierOpt.cpp @@ -0,0 +1,60 @@ +//==------------------- ExtOneapiBarrierOpt.cpp ----------------------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include +#include +#include + +using namespace sycl; + +inline thread_local uint32_t NumEventsInWaitList; + +static ur_result_t redefinedEnqueueEventsWaitWithBarrierExt(void *pParams) { + auto params = + *static_cast(pParams); + NumEventsInWaitList = *(params.pnumEventsInWaitList); + return UR_RESULT_SUCCESS; +} + +class ExtOneapiBarrierOptTest : public ::testing::Test { +public: + ExtOneapiBarrierOptTest() : Mock{} {} + +protected: + void SetUp() override { NumEventsInWaitList = 0; } + +protected: + sycl::unittest::UrMock<> Mock; +}; + +// Check that ext_oneapi_submit_barrier works fine in the scenarios +// when provided waitlist consists of only empty events. +// Tets for https://github.com/intel/llvm/pull/12951 +TEST(ExtOneapiBarrierOptTest, EmptyEventTest) { + sycl::queue q1{{sycl::property::queue::in_order()}}; + + mock::getCallbacks().set_after_callback( + "urEnqueueEventsWaitWithBarrierExt", + &redefinedEnqueueEventsWaitWithBarrierExt); + + NumEventsInWaitList = 100; + q1.ext_oneapi_submit_barrier(); + ASSERT_EQ(0u, NumEventsInWaitList); + + // ext_oneapi_submit_barrier should ignore empty, default constructed events. + sycl::event E1{}; + NumEventsInWaitList = 100; + q1.ext_oneapi_submit_barrier({E1}); + ASSERT_EQ(0u, NumEventsInWaitList); + + sycl::event E2{}; + NumEventsInWaitList = 100; + q1.ext_oneapi_submit_barrier({E1, E2}); + ASSERT_EQ(0u, NumEventsInWaitList); +} From 0060d9adec52d8a2bfc2e11b9817c860e1c96ba4 Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Wed, 24 Sep 2025 06:36:49 +0200 Subject: [PATCH 5/5] Disable in_order_ext_oneapi_submit_barrier test in 6.3 ABI testing --- devops/compat_ci_exclude.sycl-rel-6_3 | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/devops/compat_ci_exclude.sycl-rel-6_3 b/devops/compat_ci_exclude.sycl-rel-6_3 index 3bdd8ce6ea0ce..ca5479a5420af 100644 --- a/devops/compat_ci_exclude.sycl-rel-6_3 +++ b/devops/compat_ci_exclude.sycl-rel-6_3 @@ -9,6 +9,13 @@ Matrix/joint_matrix_bf16_fill_k_cache_runtime_dim.cpp # See GSD-11097. Assert/assert_in_kernels.cpp +# https://github.com/intel/llvm/pull/20159 prevents returning last event as an +# optimization for submitting barrier to an empty IOQ. However, the test +# actually checks whether last event is returned or not, so it needs to be +# updated to match the new behavior. ext_oneapi_submit_barrier spec doesn't +# require last event to be returned, so this is not an ABI break. +InorderQueue/in_order_ext_oneapi_submit_barrier.cpp + # Likely OK, but need author to provide justification, get approval/confirmation # from someone: