Skip to content

[SYCL] Graph recording support for handler-less kernel submission path #20250

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

Merged
merged 2 commits into from
Oct 6, 2025
+68 −51
Merged

[SYCL] Graph recording support for handler-less kernel submission path #20250

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
a6330e3
[SYCL] Graph recording support for handler-less kernel submission path
slawekptak Sep 30, 2025
74b007d
Address review comment
slawekptak Oct 1, 2025
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
60 changes: 60 additions & 0 deletions sycl/source/detail/queue_impl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -420,6 +420,61 @@ queue_impl::submit_impl(const detail::type_erased_cgfo_ty &CGF,
return EventImpl;
}

EventImplPtr queue_impl::submit_command_to_graph(
ext::oneapi::experimental::detail::graph_impl &GraphImpl,
std::unique_ptr<detail::CG> CommandGroup, sycl::detail::CGType CGType,
sycl::ext::oneapi::experimental::node_type UserFacingNodeType) {
auto EventImpl = detail::event_impl::create_completed_host_event();
EventImpl->setSubmittedQueue(weak_from_this());
ext::oneapi::experimental::detail::node_impl *NodeImpl = nullptr;

// GraphImpl is read and written in this scope so we lock this graph
// with full priviledges.
ext::oneapi::experimental::detail::graph_impl::WriteLock Lock(
GraphImpl.MMutex);

ext::oneapi::experimental::node_type NodeType =
UserFacingNodeType != ext::oneapi::experimental::node_type::empty
? UserFacingNodeType
: ext::oneapi::experimental::detail::getNodeTypeFromCG(CGType);

// Create a new node in the graph representing this command-group
if (isInOrder()) {
// In-order queues create implicit linear dependencies between nodes.
// Find the last node added to the graph from this queue, so our new
// node can set it as a predecessor.
std::vector<ext::oneapi::experimental::detail::node_impl *> Deps;
if (ext::oneapi::experimental::detail::node_impl *DependentNode =
GraphImpl.getLastInorderNode(this)) {
Deps.push_back(DependentNode);
}
NodeImpl = &GraphImpl.add(NodeType, std::move(CommandGroup), Deps);

// If we are recording an in-order queue remember the new node, so it
// can be used as a dependency for any more nodes recorded from this
// queue.
GraphImpl.setLastInorderNode(*this, *NodeImpl);
} else {
ext::oneapi::experimental::detail::node_impl *LastBarrierRecordedFromQueue =
GraphImpl.getBarrierDep(weak_from_this());
std::vector<ext::oneapi::experimental::detail::node_impl *> Deps;

if (LastBarrierRecordedFromQueue) {
Deps.push_back(LastBarrierRecordedFromQueue);
}
NodeImpl = &GraphImpl.add(NodeType, std::move(CommandGroup), Deps);

if (NodeImpl->MCGType == sycl::detail::CGType::Barrier) {
GraphImpl.setBarrierDep(weak_from_this(), *NodeImpl);
}
}

// Associate an event with this new node and return the event.
GraphImpl.addEventForNode(EventImpl, *NodeImpl);

return EventImpl;
}

detail::EventImplPtr queue_impl::submit_kernel_direct_impl(
const NDRDescT &NDRDesc,
std::shared_ptr<detail::HostKernelBase> &HostKernel,
Expand Down Expand Up @@ -454,6 +509,11 @@ detail::EventImplPtr queue_impl::submit_kernel_direct_impl(
CodeLoc));
CommandGroup->MIsTopCodeLoc = IsTopCodeLoc;

if (auto GraphImpl = getCommandGraph(); GraphImpl) {
Copy link
Contributor

@sergey-semenov sergey-semenov Oct 6, 2025 •
edited
Loading

Choose a reason for hiding this comment

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

Suggested change
if (auto GraphImpl = getCommandGraph(); GraphImpl) {
if (auto GraphImpl = getCommandGraph()) {

Just a non-blocking nitpick, might just be my personal preference though.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

Thanks, I will address this in a future PR.

return submit_command_to_graph(*GraphImpl, std::move(CommandGroup),
detail::CGType::Kernel);
}

return detail::Scheduler::getInstance().addCG(std::move(CommandGroup),
*this, true);
};
Expand Down
6 changes: 6 additions & 0 deletions sycl/source/detail/queue_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -626,6 +626,12 @@ class queue_impl : public std::enable_shared_from_this<queue_impl> {

bool hasCommandGraph() const { return !MGraph.expired(); }

EventImplPtr submit_command_to_graph(
ext::oneapi::experimental::detail::graph_impl &GraphImpl,
std::unique_ptr<detail::CG> CommandGroup, sycl::detail::CGType CGType,
sycl::ext::oneapi::experimental::node_type UserFacingNodeType =
ext::oneapi::experimental::node_type::empty);

unsigned long long getQueueID() { return MQueueID; }

void *getTraceEvent() { return MTraceEvent; }
Expand Down
50 changes: 2 additions & 48 deletions sycl/source/handler.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -955,54 +955,8 @@ event handler::finalize() {
// If the queue has an associated graph then we need to take the CG and pass
// it to the graph to create a node, rather than submit it to the scheduler.
if (auto GraphImpl = Queue->getCommandGraph(); GraphImpl) {
auto EventImpl = detail::event_impl::create_completed_host_event();
EventImpl->setSubmittedQueue(Queue->weak_from_this());
ext::oneapi::experimental::detail::node_impl *NodeImpl = nullptr;

// GraphImpl is read and written in this scope so we lock this graph
// with full priviledges.
ext::oneapi::experimental::detail::graph_impl::WriteLock Lock(
GraphImpl->MMutex);

ext::oneapi::experimental::node_type NodeType =
impl->MUserFacingNodeType != ext::oneapi::experimental::node_type::empty
? impl->MUserFacingNodeType
: ext::oneapi::experimental::detail::getNodeTypeFromCG(getType());

// Create a new node in the graph representing this command-group
if (Queue->isInOrder()) {
// In-order queues create implicit linear dependencies between nodes.
// Find the last node added to the graph from this queue, so our new
// node can set it as a predecessor.
std::vector<ext::oneapi::experimental::detail::node_impl *> Deps;
if (ext::oneapi::experimental::detail::node_impl *DependentNode =
GraphImpl->getLastInorderNode(Queue)) {
Deps.push_back(DependentNode);
}
NodeImpl = &GraphImpl->add(NodeType, std::move(CommandGroup), Deps);

// If we are recording an in-order queue remember the new node, so it
// can be used as a dependency for any more nodes recorded from this
// queue.
GraphImpl->setLastInorderNode(*Queue, *NodeImpl);
} else {
ext::oneapi::experimental::detail::node_impl
*LastBarrierRecordedFromQueue =
GraphImpl->getBarrierDep(Queue->weak_from_this());
std::vector<ext::oneapi::experimental::detail::node_impl *> Deps;

if (LastBarrierRecordedFromQueue) {
Deps.push_back(LastBarrierRecordedFromQueue);
}
NodeImpl = &GraphImpl->add(NodeType, std::move(CommandGroup), Deps);

if (NodeImpl->MCGType == sycl::detail::CGType::Barrier) {
GraphImpl->setBarrierDep(Queue->weak_from_this(), *NodeImpl);
}
}

// Associate an event with this new node and return the event.
GraphImpl->addEventForNode(EventImpl, *NodeImpl);
auto EventImpl = Queue->submit_command_to_graph(
*GraphImpl, std::move(CommandGroup), type, impl->MUserFacingNodeType);

#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
return EventImpl;
Expand Down
3 changes: 0 additions & 3 deletions sycl/unittests/Extensions/CommandGraph/CommandGraph.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -626,8 +626,6 @@ TEST_F(CommandGraphTest, AccessorModeEdges) {

// Tests the transitive queue recording behaviour with queue shortcuts.
TEST_F(CommandGraphTest, TransitiveRecordingShortcuts) {
// Graphs not supported yet for the no-handler submit path
#ifndef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
device Dev;
context Ctx{{Dev}};
queue Q1{Ctx, Dev};
Expand Down Expand Up @@ -671,7 +669,6 @@ TEST_F(CommandGraphTest, TransitiveRecordingShortcuts) {
ext::oneapi::experimental::queue_state::executing);
ASSERT_EQ(Q3.ext_oneapi_get_state(),
ext::oneapi::experimental::queue_state::executing);
#endif
}

// Tests that dynamic_work_group_memory.get() will throw on the host side.
Expand Down