diff --git a/sycl/include/sycl/detail/cg_types.hpp b/sycl/include/sycl/detail/cg_types.hpp index cf317fb895ba4..5e714b47cedb6 100644 --- a/sycl/include/sycl/detail/cg_types.hpp +++ b/sycl/include/sycl/detail/cg_types.hpp @@ -51,6 +51,7 @@ enum class CGType : unsigned int { EnqueueNativeCommand = 26, AsyncAlloc = 27, AsyncFree = 28, + NativeHostTask = 29, }; template struct check_fn_signature { diff --git a/sycl/source/detail/graph/node_impl.hpp b/sycl/source/detail/graph/node_impl.hpp index 3aad32f6ba601..dca597d5fb854 100644 --- a/sycl/source/detail/graph/node_impl.hpp +++ b/sycl/source/detail/graph/node_impl.hpp @@ -257,7 +257,8 @@ class node_impl : public std::enable_shared_from_this { case sycl::detail::CGType::Memset2DUSM: return createCGCopy(); case sycl::detail::CGType::EnqueueNativeCommand: - case sycl::detail::CGType::CodeplayHostTask: { + case sycl::detail::CGType::CodeplayHostTask: + case sycl::detail::CGType::NativeHostTask: { // The unique_ptr to the `sycl::detail::HostTask`, which is also used for // a EnqueueNativeCommand command, in the HostTask CG prevents from // copying the CG. We overcome this restriction by creating a new CG with diff --git a/sycl/source/detail/host_task.hpp b/sycl/source/detail/host_task.hpp index e340606b14d0f..fcd37658584c8 100644 --- a/sycl/source/detail/host_task.hpp +++ b/sycl/source/detail/host_task.hpp @@ -24,32 +24,17 @@ inline namespace _V1 { class interop_handle; namespace detail { class HostTask { - enum class HostTaskOrigin { - SYCLCoreAPI, - ExtEnqueueFunctionsAPI, - }; - std::function MHostTask; std::function MInteropTask; - HostTaskOrigin MOrigin; public: - HostTask() : MHostTask([]() {}), MOrigin(HostTaskOrigin::SYCLCoreAPI) {} - HostTask(std::function &&Func, - bool IsFromExtEnqueueFunctionsAPI = false) - : MHostTask(std::move(Func)), - MOrigin(IsFromExtEnqueueFunctionsAPI - ? HostTaskOrigin::ExtEnqueueFunctionsAPI - : HostTaskOrigin::SYCLCoreAPI) {} + HostTask() : MHostTask([]() {}) {} + HostTask(std::function &&Func) : MHostTask(std::move(Func)) {} HostTask(std::function &&Func) - : MInteropTask(std::move(Func)), MOrigin(HostTaskOrigin::SYCLCoreAPI) {} + : MInteropTask(std::move(Func)) {} bool isInteropTask() const { return !!MInteropTask; } - bool isCreatedFromEnqueueFunction() const { - return MOrigin == HostTaskOrigin::ExtEnqueueFunctionsAPI; - } - void call(HostProfilingInfo *HPI) { if (!GlobalHandler::instance().isOkToDefer()) { return; diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index c85608faa2dd1..5b667eb5de0d2 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -354,7 +354,6 @@ class DispatchHostTask { } try { - auto &Queue = HostTask.MQueue; // we're ready to call the user-defined lambda now if (HostTask.MHostTask->isInteropTask()) { assert(HostTask.MQueue && @@ -362,6 +361,7 @@ class DispatchHostTask { interop_handle IH{MReqToMem, HostTask.MQueue}; // TODO: should all the backends that support this entry point use this // for host task? + auto &Queue = HostTask.MQueue; bool NativeCommandSupport = false; Queue->getAdapter().call( detail::getSyclObjImpl(Queue->get_device())->getHandleRef(), @@ -386,47 +386,7 @@ class DispatchHostTask { IH); } } else { - if (HostTask.MHostTask->isCreatedFromEnqueueFunction()) { - bool NativeHostTaskSupport = false; - Queue->getAdapter().call( - detail::getSyclObjImpl(Queue->get_device())->getHandleRef(), - UR_DEVICE_INFO_ENQUEUE_HOST_TASK_SUPPORT_EXP, - sizeof(NativeHostTaskSupport), &NativeHostTaskSupport, nullptr); - if (NativeHostTaskSupport) { - auto NativeHostTaskData = std::make_unique( - std::move(HostTask.MHostTask->MHostTask)); - ur_event_handle_t HostTaskEvent{}; - Queue->getAdapter().call( - Queue->getHandleRef(), NativeHostTask, NativeHostTaskData.get(), - nullptr, 0, nullptr, &HostTaskEvent); - // Ownership is transferred to NativeHostTask callback on success. - (void)NativeHostTaskData.release(); - - // Wait for the host task to complete asynchronously. Since - // urEnqueueHostTaskExp executes the callback asynchronously when - // UR host task support is available, we must wait for the returned - // event before notifying completion. This ensures proper dependency - // ordering and allows profiling/async-exception handlers to see the - // actual task completion rather than the enqueue time. - if (HostTaskEvent) { - try { - Queue->getAdapter().call( - 1, &HostTaskEvent); - } catch (...) { - auto CurrentException = std::current_exception(); - Queue->getAdapter().call( - HostTaskEvent); - throw; - } - Queue->getAdapter().call( - HostTaskEvent); - } - } else { - HostTask.MHostTask->call(MThisCmd->MEvent->getHostProfilingInfo()); - } - } else { - HostTask.MHostTask->call(MThisCmd->MEvent->getHostProfilingInfo()); - } + HostTask.MHostTask->call(MThisCmd->MEvent->getHostProfilingInfo()); } } catch (...) { auto CurrentException = std::current_exception(); @@ -558,6 +518,29 @@ Command::Command( #endif } +Command::Command( + CommandType Type, queue_impl *Queue, EventImplPtr Event, + ur_exp_command_buffer_handle_t CommandBuffer, + const std::vector &SyncPoints) + : MQueue(Queue ? Queue->shared_from_this() : nullptr), + MEvent(std::move(Event)), + MPreparedDepsEvents(MEvent->getPreparedDepsEvents()), + MPreparedHostDepsEvents(MEvent->getPreparedHostDepsEvents()), MType(Type), + MCommandBuffer(CommandBuffer), MSyncPointDeps(SyncPoints) { + MWorkerQueue = MQueue; + MEnqueueStatus = EnqueueResultT::SyclEnqueueReady; + + MEvent->setCommand(this); + +#ifdef XPTI_ENABLE_INSTRUMENTATION + if (!xptiTraceEnabled()) + return; + // Obtain the stream ID so all commands can emit traces to that stream; + // copying it to the member variable to avoid ABI breakage + MStreamID = getActiveXPTIStreamID(); +#endif +} + void Command::emitInstrumentationDataProxy() { #ifdef XPTI_ENABLE_INSTRUMENTATION emitInstrumentationData(); @@ -1951,6 +1934,8 @@ static std::string_view cgTypeToString(detail::CGType Type) { return "semaphore wait"; case detail::CGType::SemaphoreSignal: return "semaphore signal"; + case detail::CGType::NativeHostTask: + return "native host task"; default: return "unknown"; break; @@ -1961,23 +1946,48 @@ ExecCGCommand::ExecCGCommand( std::unique_ptr CommandGroup, queue_impl *Queue, bool EventNeeded, ur_exp_command_buffer_handle_t CommandBuffer, const std::vector &Dependencies) - : Command(CommandType::RUN_CG, Queue, CommandBuffer, Dependencies), + : Command(CommandType::RUN_CG, Queue, makeEvent(*CommandGroup, Queue), + CommandBuffer, Dependencies), MEventNeeded(EventNeeded), MCommandGroup(std::move(CommandGroup)) { - if (MCommandGroup->getType() == detail::CGType::CodeplayHostTask) { - queue_impl *SubmitQueue = - static_cast(MCommandGroup.get())->MQueue.get(); + emitInstrumentationDataProxy(); +} + +EventImplPtr ExecCGCommand::makeEvent(const detail::CG &CG, queue_impl *Queue) { + EventImplPtr ResEvent; + + if (CG.getType() == CGType::NativeHostTask) { + const auto &HT = static_cast(CG); + ResEvent = event_impl::create_device_event(*HT.MQueue); + ResEvent->setWorkerQueue(HT.MQueue); + ResEvent->setSubmittedQueue(HT.MQueue.get()); + ResEvent->setContextImpl(HT.MQueue->getContextImpl()); + } else if (CG.getType() == CGType::CodeplayHostTask) { + const auto &HT = static_cast(CG); + ResEvent = event_impl::create_incomplete_host_event(); + queue_impl *SubmitQueue = HT.MQueue.get(); assert(SubmitQueue && "Host task command group must have a valid submit queue"); - - MEvent->setSubmittedQueue(SubmitQueue); + ResEvent->setSubmittedQueue(SubmitQueue); // Initialize host profiling info if the queue has profiling enabled. if (SubmitQueue->MIsProfilingEnabled) - MEvent->initHostProfilingInfo(); + ResEvent->initHostProfilingInfo(); + } else { + ResEvent = Queue ? event_impl::create_device_event(*Queue) + : event_impl::create_incomplete_host_event(); + if (Queue) { + ResEvent->setWorkerQueue(Queue->shared_from_this()); + ResEvent->setSubmittedQueue(Queue); + ResEvent->setContextImpl(Queue->getContextImpl()); + + if (CG.getType() == CGType::ProfilingTag) { + ResEvent->markAsProfilingTagEvent(); + } + } } - if (MCommandGroup->getType() == detail::CGType::ProfilingTag) - MEvent->markAsProfilingTagEvent(); - emitInstrumentationDataProxy(); + ResEvent->setStateIncomplete(); + + return ResEvent; } #ifdef XPTI_ENABLE_INSTRUMENTATION @@ -3499,6 +3509,36 @@ ur_result_t ExecCGCommand::enqueueImpQueue() { return UR_RESULT_SUCCESS; } + case CGType::NativeHostTask: { + CGHostTask *HostTask = static_cast(MCommandGroup.get()); + + for (ArgDesc &Arg : HostTask->MArgs) { + switch (Arg.MType) { + case kernel_param_kind_t::kind_accessor: { + Requirement *Req = static_cast(Arg.MPtr); + AllocaCommandBase *AllocaCmd = getAllocaForReq(Req); + + if (AllocaCmd) + Req->MData = AllocaCmd->getMemAllocation(); + break; + } + default: + throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), + "Unsupported arg type " + + codeToString(UR_RESULT_ERROR_INVALID_VALUE)); + } + } + + auto &Queue = HostTask->MQueue; + auto NativeHostTaskData = std::make_unique( + std::move(HostTask->MHostTask->MHostTask)); + Queue->getAdapter().call( + Queue->getHandleRef(), NativeHostTask, NativeHostTaskData.get(), + nullptr, RawEvents.size(), RawEvents.data(), Event); + (void)NativeHostTaskData.release(); + SetEventHandleOrDiscard(); + return UR_RESULT_SUCCESS; + } case CGType::EnqueueNativeCommand: { CGHostTask *HostTask = static_cast(MCommandGroup.get()); diff --git a/sycl/source/detail/scheduler/commands.hpp b/sycl/source/detail/scheduler/commands.hpp index b6f91295c08b2..4c323a51102bd 100644 --- a/sycl/source/detail/scheduler/commands.hpp +++ b/sycl/source/detail/scheduler/commands.hpp @@ -128,6 +128,13 @@ class Command { ur_exp_command_buffer_handle_t CommandBuffer = nullptr, const std::vector &SyncPoints = {}); + /// This constructor takes a pre-built event and can be used for commands, + /// which require custom logic related to how the event is created. + Command( + CommandType Type, queue_impl *Queue, EventImplPtr Event, + ur_exp_command_buffer_handle_t CommandBuffer = nullptr, + const std::vector &SyncPoints = {}); + /// \param NewDep dependency to be added /// \param ToCleanUp container for commands that can be cleaned up. /// \return an optional connection cmd to enqueue @@ -666,6 +673,8 @@ class ExecCGCommand : public Command { AllocaCommandBase *getAllocaForReq(Requirement *Req); + static EventImplPtr makeEvent(const detail::CG &CG, queue_impl *Queue); + std::unique_ptr MCommandGroup; friend class Command; diff --git a/sycl/source/detail/scheduler/scheduler.cpp b/sycl/source/detail/scheduler/scheduler.cpp index 75b1a3c7e605a..2254b5861dd5b 100644 --- a/sycl/source/detail/scheduler/scheduler.cpp +++ b/sycl/source/detail/scheduler/scheduler.cpp @@ -122,7 +122,8 @@ EventImplPtr Scheduler::addCG( NewCmd = MGraphBuilder.addCGUpdateHost(std::move(CommandGroup), AuxiliaryCmds); break; - case CGType::CodeplayHostTask: { + case CGType::CodeplayHostTask: + case CGType::NativeHostTask: { NewCmd = MGraphBuilder.addCG(std::move(CommandGroup), nullptr, AuxiliaryCmds, EventNeeded); break; @@ -199,8 +200,10 @@ void Scheduler::enqueueCommandForCG(event_impl &Event, sycl::exception( sycl::make_error_code(errc::runtime), std::string("Enqueue process failed.\n") + - __SYCL_UR_ERROR_REPORT( - NewCmd->getWorkerContext()->getBackend()) + + (NewCmd->getWorkerContext() + ? __SYCL_UR_ERROR_REPORT( + NewCmd->getWorkerContext()->getBackend()) + : " ") + sycl::detail::codeToString(Res.MErrCode)), Res.MErrCode); } diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 2e00cea153ccc..33219d3b0f4fa 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -638,7 +638,8 @@ detail::EventImplPtr handler::finalize() { std::move(impl->CGData), MCodeLoc)); break; case detail::CGType::EnqueueNativeCommand: - case detail::CGType::CodeplayHostTask: { + case detail::CGType::CodeplayHostTask: + case detail::CGType::NativeHostTask: { detail::context_impl &Context = impl->get_context(); detail::queue_impl *Queue = impl->get_queue_or_null(); CommandGroup.reset( @@ -764,7 +765,9 @@ detail::EventImplPtr handler::finalize() { assert(Queue); // Native graph recording limitation - if (type == detail::CGType::CodeplayHostTask && Queue->isNativeRecording()) { + if ((type == detail::CGType::CodeplayHostTask || + type == detail::CGType::NativeHostTask) && + Queue->isNativeRecording()) { throw sycl::exception( make_error_code(errc::feature_not_supported), "SYCL host_task is not supported in native recording mode. Use " @@ -1685,9 +1688,23 @@ void handler::SetHostTask(std::function Func) { void handler::SetHostTaskFromExtEnqueueFunctions(std::function Func) { range<1> r(1); setNDRangeDescriptor(detail::nd_range_view(r)); - impl->MHostTask.reset( - new detail::HostTask(std::move(Func), /*IsFromExtEnqueueFunctionsAPI=*/ - true)); + impl->MHostTask.reset(new detail::HostTask(std::move(Func))); + + detail::queue_impl *Queue = impl->get_queue_or_null(); + if (Queue) { + detail::adapter_impl &Adapter = getContextImpl().getAdapter(); + ur_bool_t NativeHostTaskSupport = false; + + const ur_result_t Res = Adapter.call_nocheck( + detail::getSyclObjImpl(Queue->get_device())->getHandleRef(), + UR_DEVICE_INFO_ENQUEUE_HOST_TASK_SUPPORT_EXP, + sizeof(NativeHostTaskSupport), &NativeHostTaskSupport, nullptr); + if (Res == UR_RESULT_SUCCESS && NativeHostTaskSupport) { + setType(detail::CGType::NativeHostTask); + return; + } + } + setType(detail::CGType::CodeplayHostTask); } diff --git a/sycl/test-e2e/Graph/RecordReplay/NativeRecording/exception_host_task.cpp b/sycl/test-e2e/Graph/RecordReplay/NativeRecording/exception_host_task.cpp index 88b555c51f463..d9866bf48bce2 100644 --- a/sycl/test-e2e/Graph/RecordReplay/NativeRecording/exception_host_task.cpp +++ b/sycl/test-e2e/Graph/RecordReplay/NativeRecording/exception_host_task.cpp @@ -7,8 +7,11 @@ #include "../../graph_common.hpp" +#include #include +namespace syclex = sycl::ext::oneapi::experimental; + int main() { queue Queue{property::queue::in_order{}}; @@ -40,6 +43,22 @@ int main() { return 1; } + // Try to record a native host_task - this should throw an exception + if (!expectException( + [&]() { + syclex::host_task(Queue, [=]() { + // This host task should not execute in native recording mode + for (size_t i = 0; i < N; i++) { + Data[i] = i + 100; + } + }); + }, + "host_task in native recording", sycl::errc::feature_not_supported)) { + Graph.end_recording(); + free(Data, Queue); + return 1; + } + Graph.end_recording(); free(Data, Queue); diff --git a/sycl/unittests/Extensions/EnqueueFunctionsEvents.cpp b/sycl/unittests/Extensions/EnqueueFunctionsEvents.cpp index 9e6366ce16abf..a5afddbaa6737 100644 --- a/sycl/unittests/Extensions/EnqueueFunctionsEvents.cpp +++ b/sycl/unittests/Extensions/EnqueueFunctionsEvents.cpp @@ -33,6 +33,7 @@ class EnqueueFunctionsEventsTests : public ::testing::Test { counter_urUSMEnqueuePrefetch = 0; counter_urUSMEnqueueMemAdvise = 0; counter_urEnqueueEventsWaitWithBarrier = 0; + counter_urEnqueueHostTaskExp = 0; } unittest::UrMock<> Mock; @@ -407,4 +408,26 @@ TEST_F(EnqueueFunctionsEventsTests, BarrierBeforeHostTask) { ASSERT_TRUE(HostTaskTimestamp > timestamp_urEnqueueEventsWaitWithBarrier); } +TEST_F(EnqueueFunctionsEventsTests, SubmitHostTaskNative) { + mock::getCallbacks().set_replace_callback("urEnqueueHostTaskExp", + redefined_urEnqueueHostTaskExp); + mock::getCallbacks().set_replace_callback("urDeviceGetInfo", + redefined_urDeviceGetInfo); + + oneapiext::host_task(Q, [=] {}); + + ASSERT_EQ(counter_urEnqueueHostTaskExp, size_t{1}); +} + +TEST_F(EnqueueFunctionsEventsTests, SubmitHostTaskThread) { + mock::getCallbacks().set_replace_callback("urEnqueueHostTaskExp", + redefined_urEnqueueHostTaskExp); + mock::getCallbacks().set_replace_callback("urDeviceGetInfo", + redefined_urDeviceGetInfo); + + oneapiext::host_task(Q, [=] {}); + + ASSERT_EQ(counter_urEnqueueHostTaskExp, size_t{0}); +} + } // namespace diff --git a/sycl/unittests/Extensions/FreeFunctionCommands/FreeFunctionEventsHelpers.hpp b/sycl/unittests/Extensions/FreeFunctionCommands/FreeFunctionEventsHelpers.hpp index 30417517815a9..9d744cc09502a 100644 --- a/sycl/unittests/Extensions/FreeFunctionCommands/FreeFunctionEventsHelpers.hpp +++ b/sycl/unittests/Extensions/FreeFunctionCommands/FreeFunctionEventsHelpers.hpp @@ -25,6 +25,19 @@ inline ur_result_t after_urKernelGetInfo(void *pParams) { return UR_RESULT_SUCCESS; } +template +inline ur_result_t redefined_urDeviceGetInfo(void *pParams) { + auto &Params = *reinterpret_cast(pParams); + if (*Params.ppropName == UR_DEVICE_INFO_ENQUEUE_HOST_TASK_SUPPORT_EXP) { + if (*Params.ppPropValue) + *reinterpret_cast(*Params.ppPropValue) = + NativeHostTasksSupported; + if (*Params.ppPropSizeRet) + **Params.ppPropSizeRet = sizeof(ur_bool_t); + } + return UR_RESULT_SUCCESS; +} + static size_t counter_urEnqueueKernelLaunchWithArgsExp = 0; inline ur_result_t redefined_urEnqueueKernelLaunchWithArgsExp(void *pParams) { ++counter_urEnqueueKernelLaunchWithArgsExp; @@ -84,4 +97,10 @@ inline ur_result_t after_urEnqueueEventsWaitWithBarrier(void *pParams) { return UR_RESULT_SUCCESS; } +static size_t counter_urEnqueueHostTaskExp = 0; +inline ur_result_t redefined_urEnqueueHostTaskExp(void *pParams) { + ++counter_urEnqueueHostTaskExp; + return UR_RESULT_SUCCESS; +} + } // namespace FreeFunctionEventsHelpers