[SYCL] Add SYCL_LAUNCH_BLOCKING to make submissions synchronous - #23009
Draft
uditagarwal97 wants to merge 2 commits into
Draft
[SYCL] Add SYCL_LAUNCH_BLOCKING to make submissions synchronous#23009uditagarwal97 wants to merge 2 commits into
SYCL_LAUNCH_BLOCKING to make submissions synchronous#23009uditagarwal97 wants to merge 2 commits into
Conversation
Adds a SYCL_LAUNCH_BLOCKING environment variable, the DPC++ counterpart of CUDA_LAUNCH_BLOCKING. When set to a non-zero value, every submission to any sycl::queue blocks until the queue has drained, which makes an error attributable to the launch that caused it. Because it is keyed off an environment variable rather than a queue property, it also covers queues the application does not create itself, such as those owned by oneDNN or oneMKL. The wait is placed at the funnels every queue operation passes through: both exits of queue_impl::submit_impl, queue_impl::submit_direct (shared by the kernel, graph and barrier direct paths), and both scheduler-bypass returns of queue_impl::submitMemOpHelper. Kernels, memory operations, host tasks, native commands and executable command graph submissions are all covered. Two cases are deliberately excluded. Barriers and markers run no user work and may depend on an event the application only signals after the submission returns, so waiting there would risk a hang rather than expose one. Command graph recording executes nothing, and wait() is illegal on a recording queue; native recording does not go through setCommandGraph(), so the context-level recording flag is checked alongside MGraph. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Contributor
There was a problem hiding this comment.
Pull request overview
Adds SYCL_LAUNCH_BLOCKING to make SYCL queue submissions synchronous for debugging.
Changes:
- Adds configuration parsing, documentation, and queue-wide waiting.
- Excludes barriers and graph recording from blocking.
- Adds unit and end-to-end coverage for submission paths and graphs.
Reviewed changes
Copilot reviewed 12 out of 12 changed files in this pull request and generated 1 comment.
Show a summary per file
| File | Description |
|---|---|
sycl/source/detail/config.def |
Registers the environment variable. |
sycl/source/detail/config.hpp |
Parses and caches its value. |
sycl/source/detail/queue_impl.hpp |
Defines conditional blocking behavior. |
sycl/source/detail/queue_impl.cpp |
Applies blocking across submission paths. |
sycl/doc/EnvironmentVariables.md |
Documents behavior and exclusions. |
sycl/unittests/queue/CMakeLists.txt |
Registers unit tests. |
sycl/unittests/queue/LaunchBlocking.cpp |
Tests queue draining and exclusions. |
sycl/test-e2e/Basic/launch_blocking.cpp |
Tests kernels and host tasks. |
sycl/test-e2e/Basic/launch_blocking_memops.cpp |
Tests memory operations. |
sycl/test-e2e/Basic/launch_blocking_barrier.cpp |
Tests barrier behavior. |
sycl/test-e2e/Graph/launch_blocking.cpp |
Tests command graphs. |
sycl/test-e2e/Graph/RecordReplay/NativeRecording/launch_blocking.cpp |
Tests native graph recording. |
Suppressed comments (1)
sycl/source/detail/queue_impl.hpp:479
- This context-wide check disables launch blocking for unrelated queues. If one queue starts native recording,
isNativeRecordingActive()becomes true for the entire context, so a submission to any other queue in that context returns asynchronously. Check whether this queue itself is being captured instead.
// Nothing is executed while a command graph is being recorded, and wait()
// is not a legal operation on a recording queue. Native recording does not
// go through setCommandGraph(), so MGraph is not set for it and the
// context-level flag has to be checked as well.
if (!MGraph.expired() || getContextImpl().isNativeRecordingActive())
return;
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
Comment on lines
+468
to
+469
| if (Type == CGType::Barrier || Type == CGType::BarrierWaitlist) | ||
| return; |
Three problems, all in the new E2E tests: - ext_oneapi_empty() is not a usable oracle for "the submission waited". On the OpenCL backend the query enqueues a fresh marker and reads its status without flushing, so it reports the queue as non-empty even right after the queue has been drained. The tests now check the submission's result instead: the destination is a shared allocation and the value written by the kernel or memory operation has to be visible to the host by the time the submission returns. The device-side spin loop reads its iteration count from memory so that it cannot be folded away. Verified in both directions on PVC: the checks pass in blocking mode and fail when the runtime does not block. prefetch, mem_advise and exp::prefetch are dropped along with the old oracle, as they have no observable effect on memory and take the same two exits of the fast path as the operations that are still covered. - std::getenv is deprecated on Windows, which fails the -Werror build. The tests are told about blocking mode through a command line argument from the RUN line instead, so no environment variable has to be read at all. - Graph/launch_blocking.cpp is directly under Graph/, where no lit.local.cfg requires graph support, so it ran on the OpenCL and native CPU configurations and threw. It now requires aspect-ext_oneapi_limited_graph itself. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Adds a
SYCL_LAUNCH_BLOCKINGenvironment variable, the DPC++ counterpart ofCUDA_LAUNCH_BLOCKING. When set to a non-zero value, every submission to any sycl::queue blocks until the command execution finishes.