Skip to content

[SYCL] Native host task submission without thread pool. #22518

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Draft
slawekptak wants to merge 7 commits into
base: sycl
Choose a base branch
from
Draft

[SYCL] Native host task submission without thread pool. #22518

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
643e97e
[SYCL] Native host task submission without thread pool.
slawekptak Jul 2, 2026
da2129c
New native host task command - WIP
slawekptak Jul 2, 2026
f74afae
Native host task support - WIP
slawekptak Jul 7, 2026
b420295
Pass the wait event list to the native host task UR call.
slawekptak Jul 7, 2026
e8da167
A new function to create an event for the use in ExecCGCommand,
slawekptak Jul 7, 2026
061df77
Add a restriction for native graph recording mode.
slawekptak Jul 8, 2026
7c67b4d
Move the event-related logic from ExecCGCommand constructor to makeEv…
slawekptak Jul 8, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions sycl/include/sycl/detail/cg_types.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -51,6 +51,7 @@ enum class CGType : unsigned int {
EnqueueNativeCommand = 26,
AsyncAlloc = 27,
AsyncFree = 28,
NativeHostTask = 29,
};

template <typename, typename T> struct check_fn_signature {
Expand Down
1 change: 1 addition & 0 deletions sycl/source/detail/graph/node_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -311,6 +311,7 @@ class node_impl : public std::enable_shared_from_this<node_impl> {
return createCGCopy<sycl::detail::CGAsyncAlloc>();
case sycl::detail::CGType::AsyncFree:
return createCGCopy<sycl::detail::CGAsyncFree>();
case sycl::detail::CGType::NativeHostTask:
case sycl::detail::CGType::None:
Comment on lines +314 to 315

@iclsrc iclsrc Jul 8, 2026

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🔵 Suggestion: Document the intentional fallthrough

NativeHostTask falls through silently to None and returns nullptr. This is correct — native host tasks cannot be recorded in a graph node (the exception is thrown earlier in handler::finalize) — but the silent fallthrough may confuse future readers. A brief [[fallthrough]] with a comment (or a separate return nullptr with an explanation) would make the intent clear:

Suggested change
case sycl::detail::CGType::NativeHostTask:
case sycl::detail::CGType::None:
case sycl::detail::CGType::NativeHostTask:
// NativeHostTask is not recordable; graph recording throws before reaching here.
[[fallthrough]];
case sycl::detail::CGType::None:
return nullptr;

return nullptr;
}
Expand Down
22 changes: 3 additions & 19 deletions sycl/source/detail/host_task.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,32 +24,16 @@ inline namespace _V1 {
class interop_handle;
namespace detail {
class HostTask {
enum class HostTaskOrigin {
SYCLCoreAPI,
ExtEnqueueFunctionsAPI,
};

std::function<void()> MHostTask;
std::function<void(interop_handle)> MInteropTask;
HostTaskOrigin MOrigin;

public:
HostTask() : MHostTask([]() {}), MOrigin(HostTaskOrigin::SYCLCoreAPI) {}
HostTask(std::function<void()> &&Func,
bool IsFromExtEnqueueFunctionsAPI = false)
: MHostTask(std::move(Func)),
MOrigin(IsFromExtEnqueueFunctionsAPI
? HostTaskOrigin::ExtEnqueueFunctionsAPI
: HostTaskOrigin::SYCLCoreAPI) {}
HostTask(std::function<void(interop_handle)> &&Func)
: MInteropTask(std::move(Func)), MOrigin(HostTaskOrigin::SYCLCoreAPI) {}
HostTask() : MHostTask([]() {}) {}
HostTask(std::function<void()> &&Func) : MHostTask(Func) {}
HostTask(std::function<void(interop_handle)> &&Func) : MInteropTask(Func) {}
Comment on lines +31 to +33
Comment on lines +32 to +33

@iclsrc iclsrc Jul 8, 2026

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 Important: std::move missing — constructors copy instead of move

Both constructors accept rvalue references (&&) but then copy-construct the std::function members. This silently defeats the move semantics, especially costly when the captured lambda is large. The original code used std::move here.

Suggested change
HostTask(std::function<void()> &&Func) : MHostTask(Func) {}
HostTask(std::function<void(interop_handle)> &&Func) : MInteropTask(Func) {}
HostTask(std::function<void()> &&Func) : MHostTask(std::move(Func)) {}
HostTask(std::function<void(interop_handle)> &&Func) : 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;
Expand Down
143 changes: 91 additions & 52 deletions sycl/source/detail/scheduler/commands.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -354,14 +354,14 @@ class DispatchHostTask {
}

try {
auto &Queue = HostTask.MQueue;
// we're ready to call the user-defined lambda now
if (HostTask.MHostTask->isInteropTask()) {
assert(HostTask.MQueue &&
"Host task submissions should have an associated queue");
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<UrApiKind::urDeviceGetInfo>(
detail::getSyclObjImpl(Queue->get_device())->getHandleRef(),
Expand All @@ -386,47 +386,7 @@ class DispatchHostTask {
IH);
}
} else {
if (HostTask.MHostTask->isCreatedFromEnqueueFunction()) {
bool NativeHostTaskSupport = false;
Queue->getAdapter().call<UrApiKind::urDeviceGetInfo>(
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<EnqueueHostTaskData>(
std::move(HostTask.MHostTask->MHostTask));
ur_event_handle_t HostTaskEvent{};
Queue->getAdapter().call<UrApiKind::urEnqueueHostTaskExp>(
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<UrApiKind::urEventWait>(
1, &HostTaskEvent);
} catch (...) {
auto CurrentException = std::current_exception();
Queue->getAdapter().call<UrApiKind::urEventRelease>(
HostTaskEvent);
throw;
}
Queue->getAdapter().call<UrApiKind::urEventRelease>(
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();
Expand Down Expand Up @@ -558,6 +518,27 @@ Command::Command(
#endif
}

Command::Command(
CommandType Type, queue_impl *Queue, EventImplPtr Event,
ur_exp_command_buffer_handle_t CommandBuffer,
const std::vector<ur_exp_command_buffer_sync_point_t> &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;

#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();
Expand Down Expand Up @@ -1951,6 +1932,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;
Expand All @@ -1961,23 +1944,49 @@ ExecCGCommand::ExecCGCommand(
std::unique_ptr<detail::CG> CommandGroup, queue_impl *Queue,
bool EventNeeded, ur_exp_command_buffer_handle_t CommandBuffer,
const std::vector<ur_exp_command_buffer_sync_point_t> &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<detail::CGHostTask *>(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<const CGHostTask &>(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<const CGHostTask &>(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();
ResEvent->setCommand(this);

return ResEvent;
}

#ifdef XPTI_ENABLE_INSTRUMENTATION
Expand Down Expand Up @@ -3499,6 +3508,36 @@ ur_result_t ExecCGCommand::enqueueImpQueue() {

return UR_RESULT_SUCCESS;
}
case CGType::NativeHostTask: {
CGHostTask *HostTask = static_cast<CGHostTask *>(MCommandGroup.get());

for (ArgDesc &Arg : HostTask->MArgs) {
switch (Arg.MType) {
case kernel_param_kind_t::kind_accessor: {
Requirement *Req = static_cast<Requirement *>(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<EnqueueHostTaskData>(
std::move(HostTask->MHostTask->MHostTask));
Queue->getAdapter().call<UrApiKind::urEnqueueHostTaskExp>(
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<CGHostTask *>(MCommandGroup.get());

Expand Down
7 changes: 7 additions & 0 deletions sycl/source/detail/scheduler/commands.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -128,6 +128,11 @@ class Command {
ur_exp_command_buffer_handle_t CommandBuffer = nullptr,
const std::vector<ur_exp_command_buffer_sync_point_t> &SyncPoints = {});

Command(
CommandType Type, queue_impl *Queue, EventImplPtr Event,
ur_exp_command_buffer_handle_t CommandBuffer = nullptr,
const std::vector<ur_exp_command_buffer_sync_point_t> &SyncPoints = {});

/// \param NewDep dependency to be added
/// \param ToCleanUp container for commands that can be cleaned up.
/// \return an optional connection cmd to enqueue
Expand Down Expand Up @@ -666,6 +671,8 @@ class ExecCGCommand : public Command {

AllocaCommandBase *getAllocaForReq(Requirement *Req);

EventImplPtr makeEvent(const detail::CG &CG, queue_impl *Queue);

@iclsrc iclsrc Jul 8, 2026

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 Important: makeEvent should be static

makeEvent is a non-static member function called in ExecCGCommand's constructor member-initializer list as an argument to the Command base-class constructor — meaning it runs before Command is initialized. Inside, it calls ResEvent->setCommand(this) where this points to a partially-constructed object (no base class members have been written yet).

This is technically safe only because makeEvent happens to not access any ExecCGCommand or Command members, but a future change to the function could silently read uninitialized state. Marking it static makes the constraint explicit and prevents accidental access:

Suggested change
EventImplPtr makeEvent(const detail::CG &CG, queue_impl *Queue);
static EventImplPtr makeEvent(const detail::CG &CG, queue_impl *Queue);

The corresponding definition in commands.cpp should also gain the static qualifier (or remove the class scope — either is acceptable for a private static helper).


std::unique_ptr<detail::CG> MCommandGroup;

friend class Command;
Expand Down
5 changes: 4 additions & 1 deletion sycl/source/detail/scheduler/scheduler.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down Expand Up @@ -199,6 +200,8 @@ void Scheduler::enqueueCommandForCG(event_impl &Event,
sycl::exception(
sycl::make_error_code(errc::runtime),
std::string("Enqueue process failed.\n") +
// TODO native host tasks - getWorkerContext() will return
// nullptr for native host task
__SYCL_UR_ERROR_REPORT(
NewCmd->getWorkerContext()->getBackend()) +
Comment on lines +203 to 206

@iclsrc iclsrc Jul 8, 2026

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🔴 Critical: Guaranteed null dereference on NativeHostTask enqueue failure

The TODO comment on lines 203-204 explicitly documents that getWorkerContext() returns nullptr for a NativeHostTask command, yet line 206 calls ->getBackend() on that null pointer with no guard. If urEnqueueHostTaskExp ever fails (UR error, missing adapter support at runtime, etc.) this will crash.

The fix should guard the getWorkerContext() call before dereferencing:

Suggested change
// TODO native host tasks - getWorkerContext() will return
// nullptr for native host task
__SYCL_UR_ERROR_REPORT(
NewCmd->getWorkerContext()->getBackend()) +
std::string("Enqueue process failed.\n") +
(NewCmd->getWorkerContext()
? __SYCL_UR_ERROR_REPORT(
NewCmd->getWorkerContext()->getBackend())
: std::string()) +

Alternatively, give NativeHostTask commands a non-null worker context (e.g. the queue's context) so the common error path works without special-casing.

sycl::detail::codeToString(Res.MErrCode)),
Expand Down
27 changes: 22 additions & 5 deletions sycl/source/handler.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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(
Expand Down Expand Up @@ -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 "
Expand Down Expand Up @@ -1685,9 +1688,23 @@ void handler::SetHostTask(std::function<void()> Func) {
void handler::SetHostTaskFromExtEnqueueFunctions(std::function<void()> 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();
bool NativeHostTaskSupport = false;

Adapter.call<UrApiKind::urDeviceGetInfo>(
detail::getSyclObjImpl(Queue->get_device())->getHandleRef(),
UR_DEVICE_INFO_ENQUEUE_HOST_TASK_SUPPORT_EXP,
sizeof(NativeHostTaskSupport), &NativeHostTaskSupport, nullptr);
if (NativeHostTaskSupport) {
setType(detail::CGType::NativeHostTask);
return;
}
Comment on lines +1695 to +1705
}

setType(detail::CGType::CodeplayHostTask);
}

Expand Down
19 changes: 19 additions & 0 deletions sycl/test-e2e/Graph/RecordReplay/NativeRecording/exception_host_task.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7,8 +7,11 @@

#include "../../graph_common.hpp"

#include <sycl/ext/oneapi/experimental/enqueue_functions.hpp>
#include <sycl/properties/all_properties.hpp>

namespace syclex = sycl::ext::oneapi::experimental;

int main() {
queue Queue{property::queue::in_order{}};

Expand Down Expand Up @@ -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);

Expand Down
Loading
Loading