From d28df15b5184b6df7fce2282881560f7430240b5 Mon Sep 17 00:00:00 2001 From: "Ptak, Slawomir" Date: Fri, 5 Sep 2025 11:41:08 +0000 Subject: [PATCH 1/6] [SYCL] Event-less APIs synchronization with the scheduler Commands submitted to the scheduler need to unconditionally be associated with an event (for both event and event-less APIs). This is because some commands might already be scheduled and waiting for the submission, and a newly submitted command need to return an event which can be used by the in-order type queue to properly order the commands. --- sycl/source/handler.cpp | 12 +-- .../FreeFunctionEventsHelpers.hpp | 10 --- .../scheduler/InOrderQueueHostTaskDeps.cpp | 77 ++++++++++++++++--- 3 files changed, 71 insertions(+), 28 deletions(-) diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 397f01983add4..69bae9f311d94 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -984,15 +984,15 @@ event handler::finalize() { #endif } - bool DiscardEvent = !impl->MEventNeeded && Queue && - Queue->supportsDiscardingPiEvents() && - CommandGroup->getRequirements().size() == 0; - + // Regardless of whether an event has been requested, the scheduler + // needs to generate an event so the commands are properly ordered + // (for in-order queue) and synchronized with a barrier (for out-of-order + // queue) detail::EventImplPtr Event = detail::Scheduler::getInstance().addCG( - std::move(CommandGroup), *Queue, !DiscardEvent); + std::move(CommandGroup), *Queue, true); #ifdef __INTEL_PREVIEW_BREAKING_CHANGES - return DiscardEvent ? nullptr : Event; + return Event; #else return detail::createSyclObjFromImpl(Event); #endif diff --git a/sycl/unittests/Extensions/FreeFunctionCommands/FreeFunctionEventsHelpers.hpp b/sycl/unittests/Extensions/FreeFunctionCommands/FreeFunctionEventsHelpers.hpp index c45d72ea4c343..a88b49d491eac 100644 --- a/sycl/unittests/Extensions/FreeFunctionCommands/FreeFunctionEventsHelpers.hpp +++ b/sycl/unittests/Extensions/FreeFunctionCommands/FreeFunctionEventsHelpers.hpp @@ -26,40 +26,30 @@ inline ur_result_t after_urKernelGetInfo(void *pParams) { static thread_local size_t counter_urEnqueueKernelLaunch = 0; inline ur_result_t redefined_urEnqueueKernelLaunch(void *pParams) { ++counter_urEnqueueKernelLaunch; - auto params = *static_cast(pParams); - EXPECT_EQ(*params.pphEvent, nullptr); return UR_RESULT_SUCCESS; } static thread_local size_t counter_urUSMEnqueueMemcpy = 0; inline ur_result_t redefined_urUSMEnqueueMemcpy(void *pParams) { ++counter_urUSMEnqueueMemcpy; - auto params = *static_cast(pParams); - EXPECT_EQ(*params.pphEvent, nullptr); return UR_RESULT_SUCCESS; } static thread_local size_t counter_urUSMEnqueueFill = 0; inline ur_result_t redefined_urUSMEnqueueFill(void *pParams) { ++counter_urUSMEnqueueFill; - auto params = *static_cast(pParams); - EXPECT_EQ(*params.pphEvent, nullptr); return UR_RESULT_SUCCESS; } static thread_local size_t counter_urUSMEnqueuePrefetch = 0; inline ur_result_t redefined_urUSMEnqueuePrefetch(void *pParams) { ++counter_urUSMEnqueuePrefetch; - auto params = *static_cast(pParams); - EXPECT_EQ(*params.pphEvent, nullptr); return UR_RESULT_SUCCESS; } static thread_local size_t counter_urUSMEnqueueMemAdvise = 0; inline ur_result_t redefined_urUSMEnqueueMemAdvise(void *pParams) { ++counter_urUSMEnqueueMemAdvise; - auto params = *static_cast(pParams); - EXPECT_EQ(*params.pphEvent, nullptr); return UR_RESULT_SUCCESS; } diff --git a/sycl/unittests/scheduler/InOrderQueueHostTaskDeps.cpp b/sycl/unittests/scheduler/InOrderQueueHostTaskDeps.cpp index 31a2914e2c803..1a9cf2edd9090 100644 --- a/sycl/unittests/scheduler/InOrderQueueHostTaskDeps.cpp +++ b/sycl/unittests/scheduler/InOrderQueueHostTaskDeps.cpp @@ -21,6 +21,8 @@ using namespace sycl; +namespace oneapiext = ext::oneapi::experimental; + size_t GEventsWaitCounter = 0; inline ur_result_t redefinedEventsWaitWithBarrier(void *pParams) { @@ -49,19 +51,19 @@ TEST_F(SchedulerTest, InOrderQueueHostTaskDeps) { } enum class CommandType { KERNEL = 1, MEMSET = 2 }; -std::vector> ExecutedCommands; +std::vector> ExecutedCommands; inline ur_result_t customEnqueueKernelLaunch(void *pParams) { auto params = *static_cast(pParams); - ExecutedCommands.push_back( - {CommandType::KERNEL, *params.pnumEventsInWaitList}); + ExecutedCommands.push_back({CommandType::KERNEL, *params.pnumEventsInWaitList, + *params.ppGlobalWorkSize[0]}); return UR_RESULT_SUCCESS; } inline ur_result_t customEnqueueUSMFill(void *pParams) { auto params = *static_cast(pParams); ExecutedCommands.push_back( - {CommandType::MEMSET, *params.pnumEventsInWaitList}); + {CommandType::MEMSET, *params.pnumEventsInWaitList, 0}); return UR_RESULT_SUCCESS; } @@ -112,10 +114,12 @@ TEST_F(SchedulerTest, InOrderQueueCrossDeps) { InOrderQueue.wait(); ASSERT_EQ(ExecutedCommands.size(), 2u); - EXPECT_EQ(ExecutedCommands[0].first /*CommandType*/, CommandType::MEMSET); - EXPECT_EQ(ExecutedCommands[0].second /*EventsCount*/, 0u); - EXPECT_EQ(ExecutedCommands[1].first /*CommandType*/, CommandType::KERNEL); - EXPECT_EQ(ExecutedCommands[1].second /*EventsCount*/, 0u); + EXPECT_EQ(std::get<0>(ExecutedCommands[0]) /*CommandType*/, + CommandType::MEMSET); + EXPECT_EQ(std::get<1>(ExecutedCommands[0]) /*EventsCount*/, 0u); + EXPECT_EQ(std::get<0>(ExecutedCommands[1]) /*CommandType*/, + CommandType::KERNEL); + EXPECT_EQ(std::get<1>(ExecutedCommands[1]) /*EventsCount*/, 0u); } TEST_F(SchedulerTest, InOrderQueueCrossDepsShortcutFuncs) { @@ -157,8 +161,57 @@ TEST_F(SchedulerTest, InOrderQueueCrossDepsShortcutFuncs) { InOrderQueue.wait(); ASSERT_EQ(ExecutedCommands.size(), 2u); - EXPECT_EQ(ExecutedCommands[0].first /*CommandType*/, CommandType::MEMSET); - EXPECT_EQ(ExecutedCommands[0].second /*EventsCount*/, 0u); - EXPECT_EQ(ExecutedCommands[1].first /*CommandType*/, CommandType::KERNEL); - EXPECT_EQ(ExecutedCommands[1].second /*EventsCount*/, 0u); + EXPECT_EQ(std::get<0>(ExecutedCommands[0]) /*CommandType*/, + CommandType::MEMSET); + EXPECT_EQ(std::get<1>(ExecutedCommands[0]) /*EventsCount*/, 0u); + EXPECT_EQ(std::get<0>(ExecutedCommands[1]) /*CommandType*/, + CommandType::KERNEL); + EXPECT_EQ(std::get<1>(ExecutedCommands[1]) /*EventsCount*/, 0u); } + +TEST_F(SchedulerTest, InOrderQueueCrossDepsEnqueueFunctions) { + ExecutedCommands.clear(); + sycl::unittest::UrMock<> Mock; + mock::getCallbacks().set_before_callback("urEnqueueKernelLaunch", + &customEnqueueKernelLaunch); + + sycl::platform Plt = sycl::platform(); + + context Ctx{Plt}; + queue InOrderQueue{Ctx, default_selector_v, property::queue::in_order()}; + + std::mutex CvMutex; + std::condition_variable Cv; + bool ready = false; + + InOrderQueue.submit([&](sycl::handler &CGH) { + CGH.host_task([&] { + std::unique_lock lk(CvMutex); + Cv.wait(lk, [&ready] { return ready; }); + }); + }); + + oneapiext::nd_launch( + InOrderQueue, nd_range<1>{range<1>{32}, range<1>{32}}, [](nd_item<1>) {}); + + oneapiext::nd_launch( + InOrderQueue, nd_range<1>{range<1>{64}, range<1>{32}}, [](nd_item<1>) {}); + + { + std::unique_lock lk(CvMutex); + ready = true; + } + Cv.notify_one(); + + InOrderQueue.wait(); + + ASSERT_EQ(ExecutedCommands.size(), 2u); + EXPECT_EQ(std::get<0>(ExecutedCommands[0]) /*CommandType*/, + CommandType::KERNEL); + EXPECT_EQ(std::get<1>(ExecutedCommands[0]) /*EventsCount*/, 0u); + EXPECT_EQ(std::get<2>(ExecutedCommands[0]) /*GlobalWorkSize*/, 32u); + EXPECT_EQ(std::get<0>(ExecutedCommands[1]) /*CommandType*/, + CommandType::KERNEL); + EXPECT_EQ(std::get<1>(ExecutedCommands[1]) /*EventsCount*/, 0u); + EXPECT_EQ(std::get<2>(ExecutedCommands[1]) /*GlobalWorkSize*/, 64u); +} \ No newline at end of file From 38fb8d638631c5809c07287c9bc5ce39134b2d00 Mon Sep 17 00:00:00 2001 From: "Ptak, Slawomir" Date: Fri, 5 Sep 2025 12:56:38 +0000 Subject: [PATCH 2/6] Fix a double pointer in the test and new line at end of file --- sycl/unittests/scheduler/InOrderQueueHostTaskDeps.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/unittests/scheduler/InOrderQueueHostTaskDeps.cpp b/sycl/unittests/scheduler/InOrderQueueHostTaskDeps.cpp index 1a9cf2edd9090..ae4ec01d42398 100644 --- a/sycl/unittests/scheduler/InOrderQueueHostTaskDeps.cpp +++ b/sycl/unittests/scheduler/InOrderQueueHostTaskDeps.cpp @@ -56,7 +56,7 @@ std::vector> ExecutedCommands; inline ur_result_t customEnqueueKernelLaunch(void *pParams) { auto params = *static_cast(pParams); ExecutedCommands.push_back({CommandType::KERNEL, *params.pnumEventsInWaitList, - *params.ppGlobalWorkSize[0]}); + **params.ppGlobalWorkSize}); return UR_RESULT_SUCCESS; } @@ -214,4 +214,4 @@ TEST_F(SchedulerTest, InOrderQueueCrossDepsEnqueueFunctions) { CommandType::KERNEL); EXPECT_EQ(std::get<1>(ExecutedCommands[1]) /*EventsCount*/, 0u); EXPECT_EQ(std::get<2>(ExecutedCommands[1]) /*GlobalWorkSize*/, 64u); -} \ No newline at end of file +} From 28712eedaa48ae9059269f44f686c604f07d5f81 Mon Sep 17 00:00:00 2001 From: "Ptak, Slawomir" Date: Tue, 21 Oct 2025 14:43:04 +0000 Subject: [PATCH 3/6] Update the condition for event discard, add host task tracking to the tests --- sycl/source/handler.cpp | 27 +++++++++----- .../scheduler/InOrderQueueHostTaskDeps.cpp | 35 ++++++++++++------- 2 files changed, 40 insertions(+), 22 deletions(-) diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 821cbadc17c9c..183b57527de1b 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -511,7 +511,7 @@ event handler::finalize() { // TODO checking the size of the events vector and avoiding the call is more // efficient here at this point - const bool KernelFastPath = + const bool SchedulerBypassPath = (Queue && !Graph && !impl->MSubgraphNode && !Queue->hasCommandGraph() && !impl->CGData.MRequirements.size() && !MStreamStorage.size() && (impl->CGData.MEvents.size() == 0 || @@ -521,7 +521,7 @@ event handler::finalize() { // Extract arguments from the kernel lambda, if required. // Skipping this is currently limited to simple kernels on the fast path. if (type == detail::CGType::Kernel && impl->MKernelData.getKernelFuncPtr() && - (!KernelFastPath || impl->MKernelData.hasSpecialCaptures())) { + (!SchedulerBypassPath || impl->MKernelData.hasSpecialCaptures())) { impl->MKernelData.extractArgsAndReqsFromLambda(); } @@ -633,7 +633,7 @@ event handler::finalize() { } } - if (KernelFastPath) { + if (SchedulerBypassPath) { // if user does not add a new dependency to the dependency graph, i.e. // the graph is not changed, then this faster path is used to submit // kernel bypassing scheduler and avoiding CommandGroup, Command objects @@ -879,15 +879,24 @@ event handler::finalize() { #endif } - // Regardless of whether an event has been requested, the scheduler - // needs to generate an event so the commands are properly ordered - // (for in-order queue) and synchronized with a barrier (for out-of-order - // queue) + // For kernel submission, regardless of whether an event has been requested, + // the scheduler needs to generate an event so the commands are properly + // ordered (for in-order queue) and synchronized with a barrier (for + // out-of-order queue). The event can only be skipped for the scheduler bypass + // path. + // + // For commands other than kernel submission, if an event has not been + // requested, the queue supports events discarding, and the scheduler + // might have been bypassed (not supported yet), the event can be skipped. + bool DiscardEvent = + (type != detail::CGType::Kernel && SchedulerBypassPath && + !impl->MEventNeeded && Queue->supportsDiscardingPiEvents()); + detail::EventImplPtr Event = detail::Scheduler::getInstance().addCG( - std::move(CommandGroup), *Queue, true); + std::move(CommandGroup), *Queue, !DiscardEvent); #ifdef __INTEL_PREVIEW_BREAKING_CHANGES - return Event; + return DiscardEvent ? nullptr : Event; #else return detail::createSyclObjFromImpl(Event); #endif diff --git a/sycl/unittests/scheduler/InOrderQueueHostTaskDeps.cpp b/sycl/unittests/scheduler/InOrderQueueHostTaskDeps.cpp index e2f79004eb043..e160713f5cd09 100644 --- a/sycl/unittests/scheduler/InOrderQueueHostTaskDeps.cpp +++ b/sycl/unittests/scheduler/InOrderQueueHostTaskDeps.cpp @@ -92,6 +92,7 @@ TEST_F(SchedulerTest, InOrderQueueCrossDeps) { CGH.host_task([&] { std::unique_lock lk(CvMutex); Cv.wait(lk, [&ready] { return ready; }); + ExecutedCommands.push_back({CommandType::HOST_TASK, 0, 0}); }); }); @@ -113,13 +114,15 @@ TEST_F(SchedulerTest, InOrderQueueCrossDeps) { InOrderQueue.wait(); - ASSERT_EQ(ExecutedCommands.size(), 2u); + ASSERT_EQ(ExecutedCommands.size(), 3u); EXPECT_EQ(std::get<0>(ExecutedCommands[0]) /*CommandType*/, - CommandType::MEMSET); - EXPECT_EQ(std::get<1>(ExecutedCommands[0]) /*EventsCount*/, 0u); + CommandType::HOST_TASK); EXPECT_EQ(std::get<0>(ExecutedCommands[1]) /*CommandType*/, - CommandType::KERNEL); + CommandType::MEMSET); EXPECT_EQ(std::get<1>(ExecutedCommands[1]) /*EventsCount*/, 0u); + EXPECT_EQ(std::get<0>(ExecutedCommands[2]) /*CommandType*/, + CommandType::KERNEL); + EXPECT_EQ(std::get<1>(ExecutedCommands[2]) /*EventsCount*/, 0u); } TEST_F(SchedulerTest, InOrderQueueCrossDepsShortcutFuncs) { @@ -143,6 +146,7 @@ TEST_F(SchedulerTest, InOrderQueueCrossDepsShortcutFuncs) { CGH.host_task([&] { std::unique_lock lk(CvMutex); Cv.wait(lk, [&ready] { return ready; }); + ExecutedCommands.push_back({CommandType::HOST_TASK, 0, 0}); }); }); @@ -160,13 +164,15 @@ TEST_F(SchedulerTest, InOrderQueueCrossDepsShortcutFuncs) { InOrderQueue.wait(); - ASSERT_EQ(ExecutedCommands.size(), 2u); + ASSERT_EQ(ExecutedCommands.size(), 3u); EXPECT_EQ(std::get<0>(ExecutedCommands[0]) /*CommandType*/, - CommandType::MEMSET); - EXPECT_EQ(std::get<1>(ExecutedCommands[0]) /*EventsCount*/, 0u); + CommandType::HOST_TASK); EXPECT_EQ(std::get<0>(ExecutedCommands[1]) /*CommandType*/, - CommandType::KERNEL); + CommandType::MEMSET); EXPECT_EQ(std::get<1>(ExecutedCommands[1]) /*EventsCount*/, 0u); + EXPECT_EQ(std::get<0>(ExecutedCommands[2]) /*CommandType*/, + CommandType::KERNEL); + EXPECT_EQ(std::get<1>(ExecutedCommands[2]) /*EventsCount*/, 0u); } TEST_F(SchedulerTest, InOrderQueueCrossDepsShortcutFuncsParallelFor) { @@ -229,6 +235,7 @@ TEST_F(SchedulerTest, InOrderQueueCrossDepsEnqueueFunctions) { CGH.host_task([&] { std::unique_lock lk(CvMutex); Cv.wait(lk, [&ready] { return ready; }); + ExecutedCommands.push_back({CommandType::HOST_TASK, 0, 0}); }); }); @@ -246,13 +253,15 @@ TEST_F(SchedulerTest, InOrderQueueCrossDepsEnqueueFunctions) { InOrderQueue.wait(); - ASSERT_EQ(ExecutedCommands.size(), 2u); + ASSERT_EQ(ExecutedCommands.size(), 3u); EXPECT_EQ(std::get<0>(ExecutedCommands[0]) /*CommandType*/, - CommandType::KERNEL); - EXPECT_EQ(std::get<1>(ExecutedCommands[0]) /*EventsCount*/, 0u); - EXPECT_EQ(std::get<2>(ExecutedCommands[0]) /*GlobalWorkSize*/, 32u); + CommandType::HOST_TASK); EXPECT_EQ(std::get<0>(ExecutedCommands[1]) /*CommandType*/, CommandType::KERNEL); EXPECT_EQ(std::get<1>(ExecutedCommands[1]) /*EventsCount*/, 0u); - EXPECT_EQ(std::get<2>(ExecutedCommands[1]) /*GlobalWorkSize*/, 64u); + EXPECT_EQ(std::get<2>(ExecutedCommands[1]) /*GlobalWorkSize*/, 32u); + EXPECT_EQ(std::get<0>(ExecutedCommands[2]) /*CommandType*/, + CommandType::KERNEL); + EXPECT_EQ(std::get<1>(ExecutedCommands[2]) /*EventsCount*/, 0u); + EXPECT_EQ(std::get<2>(ExecutedCommands[2]) /*GlobalWorkSize*/, 64u); } From e707d67ef7c8e52b6c5f51ad3c988460dd499c62 Mon Sep 17 00:00:00 2001 From: "Ptak, Slawomir" Date: Tue, 21 Oct 2025 14:54:48 +0000 Subject: [PATCH 4/6] Fix formatting --- sycl/unittests/scheduler/InOrderQueueHostTaskDeps.cpp | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/sycl/unittests/scheduler/InOrderQueueHostTaskDeps.cpp b/sycl/unittests/scheduler/InOrderQueueHostTaskDeps.cpp index e160713f5cd09..34447488956f7 100644 --- a/sycl/unittests/scheduler/InOrderQueueHostTaskDeps.cpp +++ b/sycl/unittests/scheduler/InOrderQueueHostTaskDeps.cpp @@ -210,10 +210,12 @@ TEST_F(SchedulerTest, InOrderQueueCrossDepsShortcutFuncsParallelFor) { InOrderQueue.wait(); ASSERT_EQ(ExecutedCommands.size(), 2u); - EXPECT_EQ(std::get<0>(ExecutedCommands[0]) /*CommandType*/, CommandType::HOST_TASK); + EXPECT_EQ(std::get<0>(ExecutedCommands[0]) /*CommandType*/, + CommandType::HOST_TASK); EXPECT_EQ(std::get<1>(ExecutedCommands[0]) /*EventsCount*/, 0u); - EXPECT_EQ(std::get<0>(ExecutedCommands[1]) /*CommandType*/, CommandType::KERNEL); - EXPECT_EQ(std::get<1>(ExecutedCommands[1])/*EventsCount*/, 0u); + EXPECT_EQ(std::get<0>(ExecutedCommands[1]) /*CommandType*/, + CommandType::KERNEL); + EXPECT_EQ(std::get<1>(ExecutedCommands[1]) /*EventsCount*/, 0u); } TEST_F(SchedulerTest, InOrderQueueCrossDepsEnqueueFunctions) { From e8831aa915944731683b0eb430fa29f442849f14 Mon Sep 17 00:00:00 2001 From: "Ptak, Slawomir" Date: Mon, 3 Nov 2025 13:42:01 +0000 Subject: [PATCH 5/6] Update the tests, rename the scheduler bypass variable. --- sycl/source/handler.cpp | 8 ++++---- .../FreeFunctionCommands/FreeFunctionEventsHelpers.hpp | 8 ++++++++ 2 files changed, 12 insertions(+), 4 deletions(-) diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index e3b93d187c75e..50102355beb57 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -511,7 +511,7 @@ event handler::finalize() { // TODO checking the size of the events vector and avoiding the call is more // efficient here at this point - const bool SchedulerBypassPath = + const bool KernelSchedulerBypass = (Queue && !Graph && !impl->MSubgraphNode && !Queue->hasCommandGraph() && !impl->CGData.MRequirements.size() && !MStreamStorage.size() && (impl->CGData.MEvents.size() == 0 || @@ -521,7 +521,7 @@ event handler::finalize() { // Extract arguments from the kernel lambda, if required. // Skipping this is currently limited to simple kernels on the fast path. if (type == detail::CGType::Kernel && impl->MKernelData.getKernelFuncPtr() && - (!SchedulerBypassPath || impl->MKernelData.hasSpecialCaptures())) { + (!KernelSchedulerBypass || impl->MKernelData.hasSpecialCaptures())) { impl->MKernelData.extractArgsAndReqsFromLambda(); } @@ -633,7 +633,7 @@ event handler::finalize() { } } - if (SchedulerBypassPath) { + if (KernelSchedulerBypass) { // if user does not add a new dependency to the dependency graph, i.e. // the graph is not changed, then this faster path is used to submit // kernel bypassing scheduler and avoiding CommandGroup, Command objects @@ -889,7 +889,7 @@ event handler::finalize() { // requested, the queue supports events discarding, and the scheduler // might have been bypassed (not supported yet), the event can be skipped. bool DiscardEvent = - (type != detail::CGType::Kernel && SchedulerBypassPath && + (type != detail::CGType::Kernel && KernelSchedulerBypass && !impl->MEventNeeded && Queue->supportsDiscardingPiEvents()); detail::EventImplPtr Event = detail::Scheduler::getInstance().addCG( diff --git a/sycl/unittests/Extensions/FreeFunctionCommands/FreeFunctionEventsHelpers.hpp b/sycl/unittests/Extensions/FreeFunctionCommands/FreeFunctionEventsHelpers.hpp index cd9ba7c8dbe47..a87dda42021d1 100644 --- a/sycl/unittests/Extensions/FreeFunctionCommands/FreeFunctionEventsHelpers.hpp +++ b/sycl/unittests/Extensions/FreeFunctionCommands/FreeFunctionEventsHelpers.hpp @@ -46,24 +46,32 @@ inline ur_result_t redefined_urEnqueueKernelLaunchWithEvent(void *pParams) { static thread_local size_t counter_urUSMEnqueueMemcpy = 0; inline ur_result_t redefined_urUSMEnqueueMemcpy(void *pParams) { ++counter_urUSMEnqueueMemcpy; + auto params = *static_cast(pParams); + EXPECT_EQ(*params.pphEvent, nullptr); return UR_RESULT_SUCCESS; } static thread_local size_t counter_urUSMEnqueueFill = 0; inline ur_result_t redefined_urUSMEnqueueFill(void *pParams) { ++counter_urUSMEnqueueFill; + auto params = *static_cast(pParams); + EXPECT_EQ(*params.pphEvent, nullptr); return UR_RESULT_SUCCESS; } static thread_local size_t counter_urUSMEnqueuePrefetch = 0; inline ur_result_t redefined_urUSMEnqueuePrefetch(void *pParams) { ++counter_urUSMEnqueuePrefetch; + auto params = *static_cast(pParams); + EXPECT_EQ(*params.pphEvent, nullptr); return UR_RESULT_SUCCESS; } static thread_local size_t counter_urUSMEnqueueMemAdvise = 0; inline ur_result_t redefined_urUSMEnqueueMemAdvise(void *pParams) { ++counter_urUSMEnqueueMemAdvise; + auto params = *static_cast(pParams); + EXPECT_EQ(*params.pphEvent, nullptr); return UR_RESULT_SUCCESS; } From 7994f842a6a80773b11ca8659e50279bef26a51a Mon Sep 17 00:00:00 2001 From: Slawomir Ptak Date: Mon, 3 Nov 2025 16:32:18 +0100 Subject: [PATCH 6/6] Apply suggestion from @sergey-semenov Co-authored-by: Sergey Semenov --- sycl/source/handler.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 50102355beb57..f9440c089f2a0 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -887,7 +887,7 @@ event handler::finalize() { // // For commands other than kernel submission, if an event has not been // requested, the queue supports events discarding, and the scheduler - // might have been bypassed (not supported yet), the event can be skipped. + // could have been bypassed (not supported yet), the event can be skipped. bool DiscardEvent = (type != detail::CGType::Kernel && KernelSchedulerBypass && !impl->MEventNeeded && Queue->supportsDiscardingPiEvents());