Skip to content

[SYCL][Graph] Avoid removing in-order event dependencies across the native recording capture boundary #22604

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

Open
mmichel11 wants to merge 4 commits into
base: sycl
Choose a base branch
from
Open

[SYCL][Graph] Avoid removing in-order event dependencies across the native recording capture boundary #22604

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
df3876e
[SYCL][Graph] Avoid removing in-order event dependencies across the n…
mmichel11 Jul 10, 2026
b918068
Adjust comments
mmichel11 Jul 12, 2026
9742ceb
Address review feedback
mmichel11 Jul 16, 2026
37b6d87
Remove unneeded include and assert in destructor
mmichel11 Jul 16, 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
24 changes: 24 additions & 0 deletions sycl/source/detail/context_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -272,6 +272,24 @@ class context_impl : public std::enable_shared_from_this<context_impl> {
bool supportsReusableEvents();
bool supportsEventProfiling();

/// @return True if any native graph recording is active on a queue
/// in this context.
bool isNativeRecordingActive() const {
return MNativeRecordingCount.load(std::memory_order_acquire) > 0;

@adamfidel adamfidel Jul 15, 2026

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

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

I think we can replace the acquire/release stuff with memory_order_relaxed since I don't think a stricter ordering buys us anything.

@mmichel11 mmichel11 Jul 16, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

The reason I used acquire-release is due to the transition of the counter from 1 -> 0 in end capture. The release update and acquire load guarantee the effects of the end capture are observed by another thread when it sees 0 in the counter. I was thinking of this for the case where a submission to a forked queue is done concurrent to end capture on the primary queue. This is always a user application race though for what is actually captured.

What do you think?

@adamfidel adamfidel Jul 16, 2026

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

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

In that case, I think keeping acquire-release is fine as is.

}

/// Register that a native graph recording has begun on a queue in this
/// context.
void nativeRecordingBegan() {
MNativeRecordingCount.fetch_add(1, std::memory_order_release);
}

/// Register that a native graph recording has ended on a queue in this
/// context.
void nativeRecordingEnded() {
MNativeRecordingCount.fetch_sub(1, std::memory_order_release);

@againull againull Jul 17, 2026 •
edited
Loading

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

I think we should add an assertion that we don't get negative counter here which indicates imbalance in Began/Ended.

  void nativeRecordingEnded() {
    [[maybe_unused]] int64_t Prev =
        MNativeRecordingCount.fetch_sub(1, std::memory_order_release);
    assert(Prev > 0 && "Native recording counter underflow");
  }

}

private:
bool MOwnedByRuntime;
async_handler MAsyncHandler;
Expand Down Expand Up @@ -365,6 +383,12 @@ class context_impl : public std::enable_shared_from_this<context_impl> {
MNativeGraphRegistry;
mutable std::mutex MNativeGraphRegistryMutex;

// Count of active native graph recordings in this context. Incremented when
// a native capture begins on a queue and decremented when it ends. Used to
// preserve in-order dependencies that cross the native-recording capture
// boundary.
std::atomic<int64_t> MNativeRecordingCount{0};

void verifyProps(const property_list &Props) const;
};

Expand Down
14 changes: 14 additions & 0 deletions sycl/source/detail/event_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -346,6 +346,14 @@ class event_impl {
return MEventFromSubmittedExecCommandBuffer;
}

bool isPotentiallyNativeRecorded() const {
return MPotentiallyNativeRecorded;
}

void setPotentiallyNativeRecorded(bool Value) {
MPotentiallyNativeRecorded = Value;
}

void setProfilingEnabled(bool Value) { MIsProfilingEnabled = Value; }

// Sets a command-buffer command when this event represents an enqueue to a
Expand Down Expand Up @@ -439,6 +447,12 @@ class event_impl {
/// Indicates that the event results from a command graph submission.
bool MEventFromSubmittedExecCommandBuffer = false;

/// Set from the context of the worker queue when the event is created for a
/// command submission, marking it as potentially captured if a native graph
/// recording was active. Used to preserve in-order dependencies that cross
/// the native-recording capture boundary.
bool MPotentiallyNativeRecorded = false;

// If this event represents a submission to a
// ur_exp_command_buffer_sync_point_t the sync point for that submission is
// stored here.
Expand Down
53 changes: 20 additions & 33 deletions sycl/source/detail/graph/graph_impl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -726,23 +726,20 @@ void graph_impl::clearQueues(bool NeedsLock) {
for (auto &Queue : SwappedQueues) {
if (auto ValidQueue = Queue.lock(); ValidQueue) {
if (MNativeGraphHandle) {
// End native UR graph capture
auto UrQueue = ValidQueue->getHandleRef();
ur_exp_graph_handle_t CapturedGraph = nullptr;
context_impl &ContextImpl = *sycl::detail::getSyclObjImpl(MContext);
sycl::detail::adapter_impl &Adapter = ContextImpl.getAdapter();
ur_result_t Result = Adapter.call_nocheck<
sycl::detail::UrApiKind::urQueueEndGraphCaptureExp>(UrQueue,
&CapturedGraph);
if (Result != UR_RESULT_SUCCESS) {
auto EndResult = ValidQueue->endNativeRecording();
if (EndResult.Result != UR_RESULT_SUCCESS) {
throw sycl::exception(sycl::make_error_code(errc::runtime),
"Failed to end native graph capture");
"Error when ending native graph capture");
}
// CapturedGraph should be the same as MNativeGraphHandle
} else {
// Only call setCommandGraph for traditional recording
ValidQueue->setCommandGraph(nullptr);
}
} else if (MNativeGraphHandle) {
// The primary recording queue was destroyed by the user. We must update
// the context that the recording is over.
getContextImpl().nativeRecordingEnded();
}
}
}
Expand Down Expand Up @@ -898,19 +895,15 @@ void graph_impl::beginRecordingImpl(sycl::detail::queue_impl &Queue,

// Use native UR graph recording if enabled
if (MNativeGraphHandle) {
auto UrQueue = Queue.getHandleRef();
context_impl &ContextImpl = *sycl::detail::getSyclObjImpl(MContext);
sycl::detail::adapter_impl &Adapter = ContextImpl.getAdapter();

if (Queue.isNativeRecording()) {
throw sycl::exception(sycl::make_error_code(errc::invalid),
"Queue is already in native graph capture mode");
}

ur_result_t Result = Adapter.call_nocheck<
sycl::detail::UrApiKind::urQueueBeginCaptureIntoGraphExp>(
UrQueue, MNativeGraphHandle);
if (Result != UR_RESULT_SUCCESS) {
auto BeginResult = Queue.beginNativeRecording(MNativeGraphHandle);
if (BeginResult.RecordingActive) {
addQueue(Queue);
}
if (BeginResult.Result != UR_RESULT_SUCCESS) {
throw sycl::exception(sycl::make_error_code(errc::runtime),
"Failed to begin native UR graph capture");
}
Expand All @@ -921,8 +914,8 @@ void graph_impl::beginRecordingImpl(sycl::detail::queue_impl &Queue,
} else {
Queue.setCommandGraphUnlocked(shared_from_this());
}
addQueue(Queue);
}
addQueue(Queue);
}
}

Expand Down Expand Up @@ -2310,22 +2303,16 @@ void modifiable_command_graph::end_recording(queue &RecordingQueue) {
// End native UR graph capture
assert(impl->getNativeGraphHandle() &&
"Native graph handle must be valid when ending native recording");
auto UrQueue = QueueImpl.getHandleRef();
ur_exp_graph_handle_t CapturedGraph = nullptr;
context_impl &ContextImpl =
*sycl::detail::getSyclObjImpl(impl->getContext());
sycl::detail::adapter_impl &Adapter = ContextImpl.getAdapter();
ur_result_t Result =
Adapter
.call_nocheck<sycl::detail::UrApiKind::urQueueEndGraphCaptureExp>(
UrQueue, &CapturedGraph);
if (Result != UR_RESULT_SUCCESS) {
auto EndResult = QueueImpl.endNativeRecording();
if (!EndResult.RecordingActive) {
impl->removeQueue(QueueImpl);
}
if (EndResult.Result != UR_RESULT_SUCCESS) {
throw sycl::exception(sycl::make_error_code(errc::runtime),
"Failed to end native UR graph capture");
"Error when ending native graph capture");
}
assert(CapturedGraph == impl->getNativeGraphHandle() &&
assert(EndResult.CapturedGraph == impl->getNativeGraphHandle() &&
"Captured graph handle must match the graph's native handle");
impl->removeQueue(QueueImpl);
}
} else {
// Traditional recording path
Expand Down
32 changes: 32 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 @@ -464,6 +464,8 @@ EventImplPtr queue_impl::submit_kernel_scheduler_bypass(
EnqueueKernel();
} else {
ResultEvent->setWorkerQueue(weak_from_this());
ResultEvent->setPotentiallyNativeRecorded(
getContextImpl().isNativeRecordingActive());
ResultEvent->setStateIncomplete();
ResultEvent->setSubmissionTime();

Expand Down Expand Up @@ -514,6 +516,8 @@ EventImplPtr queue_impl::submit_barrier_scheduler_bypass(
ResEvent->setQueue(*this);
}
ResEvent->setWorkerQueue(weak_from_this());
ResEvent->setPotentiallyNativeRecorded(
getContextImpl().isNativeRecordingActive());
ResEvent->setSubmissionTime();
ResEvent->setEnqueued();
ResEvent->setStateIncomplete();
Expand Down Expand Up @@ -660,6 +664,34 @@ bool queue_impl::isNativeRecording() const {
return Result == UR_RESULT_SUCCESS && IsGraphCaptureEnabled;
}

queue_impl::NativeRecordingResult
queue_impl::beginNativeRecording(ur_exp_graph_handle_t Graph) {
std::lock_guard<std::mutex> Lock(MMutex);
NativeRecordingResult BeginResult;
BeginResult.Result =
getAdapter().call_nocheck<UrApiKind::urQueueBeginCaptureIntoGraphExp>(
MQueue, Graph);
BeginResult.RecordingActive = BeginResult.Result == UR_RESULT_SUCCESS;
if (BeginResult.RecordingActive) {
getContextImpl().nativeRecordingBegan();
}
return BeginResult;
}

queue_impl::NativeRecordingResult queue_impl::endNativeRecording() {
std::lock_guard<std::mutex> Lock(MMutex);
NativeRecordingResult EndResult;
EndResult.Result =
getAdapter().call_nocheck<UrApiKind::urQueueEndGraphCaptureExp>(
MQueue, &EndResult.CapturedGraph);
EndResult.RecordingActive =
EndResult.Result != UR_RESULT_SUCCESS && isNativeRecording();
if (!EndResult.RecordingActive) {
getContextImpl().nativeRecordingEnded();
}
return EndResult;
}

ext::oneapi::experimental::queue_state
queue_impl::ext_oneapi_get_state_impl() const {
// A graph may either be recording at the SYCL level or recording at a lower
Expand Down
10 changes: 10 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 @@ -664,6 +664,16 @@ class queue_impl : public std::enable_shared_from_this<queue_impl> {

bool isNativeRecording() const;

struct NativeRecordingResult {
ur_exp_graph_handle_t CapturedGraph = nullptr;
bool RecordingActive = false;
ur_result_t Result = UR_RESULT_SUCCESS;
};

NativeRecordingResult beginNativeRecording(ur_exp_graph_handle_t Graph);

NativeRecordingResult endNativeRecording();

ext::oneapi::experimental::queue_state ext_oneapi_get_state_impl() const;

std::shared_ptr<ext::oneapi::experimental::detail::graph_impl>
Expand Down
29 changes: 25 additions & 4 deletions sycl/source/detail/scheduler/commands.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -209,6 +209,10 @@ static std::string commandToName(Command::CommandType Type) {
std::vector<ur_event_handle_t> Command::getUrEvents(events_range Events,
queue_impl *CommandQueue,
bool IsHostTaskCommand) {
const bool CanRemoveRedundantEvent =
CommandQueue && CommandQueue->isInOrder() && !IsHostTaskCommand &&
!CommandQueue->getContextImpl().isNativeRecordingActive();

std::vector<ur_event_handle_t> RetUrEvents;
for (event_impl &Event : Events) {
auto Handle = Event.getHandle();
Expand All @@ -219,8 +223,12 @@ std::vector<ur_event_handle_t> Command::getUrEvents(events_range Events,
// At this stage dependency is definitely ur task and need to check if
// current one is a host task. In this case we should not skip ur event due
// to different sync mechanisms for different task types on in-order queue.
if (CommandQueue && Event.getWorkerQueue().get() == CommandQueue &&
CommandQueue->isInOrder() && !IsHostTaskCommand)
// When native recording is in-progress, then we must skip this optimization
// to let the driver runtime handle event dependencies crossing the graph
// capture boundary.
if (CanRemoveRedundantEvent &&
Event.getWorkerQueue().get() == CommandQueue &&
!Event.isPotentiallyNativeRecorded())
continue;

RetUrEvents.push_back(Handle);
Expand Down Expand Up @@ -532,8 +540,11 @@ Command::Command(
if (Queue)
MEvent->setSubmittedQueue(Queue);
MEvent->setCommand(this);
if (MQueue)
MEvent->setContextImpl(MQueue->getContextImpl());
if (MQueue) {
context_impl &Context = MQueue->getContextImpl();
MEvent->setContextImpl(Context);
MEvent->setPotentiallyNativeRecorded(Context.isNativeRecordingActive());
}
MEvent->setStateIncomplete();
MEnqueueStatus = EnqueueResultT::SyclEnqueueReady;

Expand Down Expand Up @@ -1512,6 +1523,11 @@ MemCpyCommand::MemCpyCommand(const Requirement &SrcReq,

MWorkerQueue = !MQueue ? MSrcQueue : MQueue;
MEvent->setWorkerQueue(MWorkerQueue);
// When MQueue is non-null the base Command constructor already set this from
// MQueue's context.
if (!MQueue && MWorkerQueue)
MEvent->setPotentiallyNativeRecorded(
MWorkerQueue->getContextImpl().isNativeRecordingActive());

emitInstrumentationDataProxy();
}
Expand Down Expand Up @@ -1689,6 +1705,11 @@ MemCpyCommandHost::MemCpyCommandHost(const Requirement &SrcReq,

MWorkerQueue = !MQueue ? MSrcQueue : MQueue;
MEvent->setWorkerQueue(MWorkerQueue);
// When MQueue is non-null the base Command constructor already set this from
// MQueue's context.
if (!MQueue && MWorkerQueue)
MEvent->setPotentiallyNativeRecorded(
MWorkerQueue->getContextImpl().isNativeRecordingActive());

emitInstrumentationDataProxy();
}
Expand Down
Loading
Loading