[SYCL][Graph] Avoid removing in-order event dependencies across the native recording capture boundary#22604
Conversation
…ative recording capture boundary
4eae1bd to
29853c5
Compare
There was a problem hiding this comment.
Pull request overview
This PR fixes incorrect pruning of redundant in-order event wait-list dependencies when SYCL graphs use native backend recording, by conservatively tracking native recording at the context level and propagating that state onto events so dependencies that may cross the native capture boundary are preserved.
Changes:
- Add a context-level “native recording active” counter and use it to disable same-queue in-order dependency pruning during native recording.
- Tag events created while native recording is active as
PotentiallyNativeRecordedand prevent pruning of those dependencies even after recording ends. - Refactor UR begin/end capture into
queue_impl::{beginNativeRecording,endNativeRecording}and add a new end-to-end test validating the wait-list behavior viaSYCL_UR_TRACE.
Reviewed changes
Copilot reviewed 7 out of 7 changed files in this pull request and generated 3 comments.
Show a summary per file
| File | Description |
|---|---|
| sycl/test-e2e/Graph/RecordReplay/NativeRecording/inorder_event_dependencies.cpp | New E2E test asserting wait-list dependency preservation across native capture boundaries. |
| sycl/source/detail/scheduler/commands.cpp | Gate redundant in-order dependency pruning on context native-recording activity and event tagging. |
| sycl/source/detail/queue_impl.hpp | Expose queue-level wrappers for UR native capture begin/end. |
| sycl/source/detail/queue_impl.cpp | Implement begin/end native capture wrappers and mark scheduler-bypass events as potentially native recorded. |
| sycl/source/detail/graph/graph_impl.cpp | Use the new queue_impl wrappers for native capture begin/end and handle queue-destruction cleanup. |
| sycl/source/detail/event_impl.hpp | Add PotentiallyNativeRecorded flag to events to preserve cross-boundary dependencies. |
| sycl/source/detail/context_impl.hpp | Add atomic native-recording counter and query helpers on context. |
| void queue_impl::beginNativeRecording(ur_exp_graph_handle_t Graph) { | ||
| std::lock_guard<std::mutex> Lock(MMutex); | ||
| ur_result_t Result = | ||
| getAdapter().call_nocheck<UrApiKind::urQueueBeginCaptureIntoGraphExp>( | ||
| MQueue, Graph); | ||
| if (Result != UR_RESULT_SUCCESS) { | ||
| throw sycl::exception(sycl::make_error_code(errc::runtime), | ||
| "Failed to begin native UR graph capture"); | ||
| } | ||
| getContextImpl().nativeRecordingBegan(); | ||
| } |
There was a problem hiding this comment.
This scenario occurs when the user has an application race between begin_recording(q) and a command submission to the same queue. It cannot occur between different queues because the initial submission must be done on the primary queue and forking the graph to a second queue requires some form of synchronization. Since we acquire the queue lock now in this call, I believe there is no risk.
Wait event dependencies across the native graph capture boundary cannot be removed and must be processed by the native backend. Due to transitive queue recording and interoperability with native backend calls (e.g. capturing direct L0 submissions), SYCL alone cannot track if a queue is recording. We can only determine if a queue is truly recording through an adapter call which is too expensive to perform on each command submission and muddles tracing tool output.
We solve this issue when a native recording is in progress while keeping the SYCL optimization for the general case with the following approach:
PotentiallyNativeRecorded.PotentiallyNativeRecordedwait event or queue state, disable the in-order queue optimization removing the dependency.