Run CUDA delegates on the per-thread stream by shoumikhin · Pull Request #22318 · pytorch/executorch

shoumikhin · 2026-08-30T06:36:04Z

Each delegate handle created its own CUDA stream. Delegates in one program run one
after another, so a stream apiece bought no ordering between them: a delegate could
read an input while the delegate that produces it still had work queued on a
different stream.

Before this change the stream was per handle, so two CUDA delegates in one program on
one thread also had two streams and were not ordered either. That is why both in-repo
callers set the shared-stream option. A program split across the CUDA and TensorRT
backends is the case that shows it worst: the TensorRT delegate runs on the per-thread
stream when no caller stream is installed, so the two backends sat on different
streams with nothing between them. Such a program returned a different answer on
almost every call: executing one loaded program forty times produced nine distinct
output sums, one of them correct.

Every handle now takes the per-thread stream, which makes the ordering the graph
expresses the ordering the device sees, in either direction and between two CUDA
delegates, with no event handshake needed because both backends end up on the same
stream. A caller stream still takes precedence per execute.

The guarantee is per thread, and that is the one thing to get right about this change.
cudaStreamPerThread resolves to a different stream on each host thread, so delegates
called from one thread are ordered and delegates called from different threads are
not.

That is a real loss for two cases, and neither is limited to callers of the old
option. use_shared_cuda_stream created a single stream held by the backend and handed
it to every handle on every thread, so a caller running the encoder on one thread and
the decoder on another got ordering from it. Separately, a handle's single stream used
to be installed by whichever thread called execute, so a caller that guards one loaded
program with a lock and drives it from a different thread each time was ordered before
and is not now.

Because of that, the option is refused rather than ignored: setting it to true returns
NotSupported, the same shape the graph option already uses for ROCm. Accepting it
silently would leave such a caller building, running, and getting wrong numbers, and
the log line alone is not enough, since a default Release build compiles ET_LOG out
entirely. Both in-repo callers drive their methods from one thread and no longer set
it.

Two smaller consequences of one stream rather than many. Two independent programs
submitted from one thread with no synchronization between them now run one after the
other rather than overlapping, measured at about 2x on two equal single-block kernels;
a caller stream per program restores the overlap, except for a method with CUDA graphs
enabled, where execute refuses a caller stream. And a handle no longer owns a stream
that destroy could free while the thread-local table still pointed at it, which was a
real dangling-stream path before.

Because one stream is now shared, a capture that an error abandons would otherwise
leave that stream capturing for every later delegate on the thread. Two guards handle
this. One ends the capture, and is declared so that it runs before the tensor cleanup:
freeing a device buffer on a still-capturing stream fails with invalid argument and
leaks the block, measured at 64 MiB not returned on an H100. The other frees the
buffers the attempt pinned and rewinds the method to warmup, so a retry captures from
a clean state. Without that second half the retry appends a second set of static
buffers, replay reads one set while the input copies target the other, and every
execute returns the values captured at that moment with nothing reporting an error.

Test plan:

backends/cuda/tests/test_coalesced_determinism.py exports a program split across both
backends, checks the saved program contains both backends, compares the first result
against eager, then runs it a hundred times on one loaded program and requires every
result to equal the first exactly. The eager comparison is what catches a delegate
that reads a stale buffer the same way every call, which is self-consistent and still
wrong.

It passes with this change and fails without it at the second run. Each step of the
model carries its own buffer: one constant read by every island collapses to the same
placeholder name several times in the flattened graph, which is a separate export
defect and would fail this test before it reached the run loop.

Verified separately on an H100, since no test covers it:

a free issued into a still-capturing stream returns invalid argument and the block
is not returned, which is why the end-capture guard runs first; with the guards in
this order the unwind ends the capture, deletes the tensors, then frees, and nothing
leaks

one executable graph replayed from two threads on two streams returns no error and
simply serializes, so the graph path does not need a new check for that case

a capture step that fails after capture began leaves the stream clean and the
method's buffers freed, and three replays after the retry return the current answer

Measured on Linux aarch64, sm_110, on a program of twenty-five delegates: before,
thirty-nine of forty runs were wrong; after, seven thousand six hundred consecutive
runs were correct across several processes. Programs of a single delegate on either
backend were already correct and stayed so over five hundred runs each. Median latency
for that split program went from about 850 to about 940 microseconds, which is the
ordering that was previously skipped. A single delegate program is unchanged.

The three existing C++ test binaries for this backend pass. The Python suite for this
backend has eighteen failures on this machine both with and without this change, so it
introduces no regressions; those are an unrelated gap in matmul and convolution
lowering on this architecture.

cudaStreamPerThread has no alias in the HIP compatibility header, so this adds one, or
the ROCm build of this backend would not compile. It is a macro rather than an inline
constexpr because the HIP definition casts an integer to a pointer type. That alias
was checked by hand against the ROCm definitions and not by a ROCm build.

Not covered: the test above needs the TensorRT delegate, which nothing in the repo
installs, so it is collected and skipped, and the job still passes. There is therefore
no executed test for this behaviour in CI. Installing the delegate there would not be
enough either, because that job runs on x86 where this reordering does not reproduce;
real coverage needs a job on an architecture where it does.

Also not covered, because this branch is on a fork and the fork guard skips them: the
ROCm jobs, the CUDA Windows jobs, the CUDA end-to-end model job, and the CUDA pybind
job that chains off it. Those last two are the ones that push real models through the
runtime this changes on a Linux GPU. And the CUDA graph capture and replay paths beyond
confirming that a program using them runs and agrees with itself over forty runs.

pytorch-bot · 2026-08-30T06:36:09Z

🔗 Helpful Links

🧪 See artifacts and rendered test results at hud.pytorch.org/pr/pytorch/executorch/22318

📄 Preview Python docs built from this PR

Note: Links to docs will display an error until the docs builds have been completed.

❌ 3 New Failures, 2 Unrelated Failures

As of commit f823796 with merge base 2b3a32d ():

NEW FAILURES - The following jobs have failed:

pull / unittest / linux / linux-job (gh)
pytorch_tokenizers
pull / unittest-buck / linux / linux-job (gh)
pytorch_tokenizers
pull / unittest-editable / linux / linux-job (gh)
pytorch_tokenizers

FLAKY - The following job failed but was likely due to flakiness present on trunk:

pull / test-qnn-testsuite-linux / test-backend-linux (qnn, models) / linux-job (gh) (detected as infra flaky with no log or failing log classifier)

BROKEN TRUNK - The following job failed but was present on the merge base:

👉 Rebase onto the `viable/strict` branch to avoid these failures

pull / test-moshi-linux / linux-job (gh) (trunk failure)
E: Failed to fetch http://security.ubuntu.com/ubuntu/pool/main/libs/libssh/libssh-gcrypt-4_0.9.6-2ubuntu0.22.04.7_amd64.deb 404 Not Found [IP: 91.189.91.82 80]

This comment was automatically generated by Dr. CI and updates every 15 minutes.

linux-foundation-easycla · 2026-08-30T06:36:11Z

✅ login: shoumikhin / name: Anthony Shoumikhin (374821d, 828c95c)
❌ The email address for the commit (b3c12b1) is not linked to the GitHub account, preventing the EasyCLA check. Consult this Help Article and GitHub Help to resolve. (To view the commit's email address, add .patch at the end of this PR page's URL.) For further assistance with EasyCLA, please visit our EasyCLA portal and chat with our support bot.

Copilot

Copilot was unable to review this pull request because the user who requested the review has reached their quota limit.

Copilot

Copilot was unable to review this pull request because the user who requested the review has reached their quota limit.

Copilot

Copilot was unable to review this pull request because the user who requested the review has reached their quota limit.

Copilot

Copilot was unable to review this pull request because the user who requested the review has reached their quota limit.

Copilot

Copilot was unable to review this pull request because the user who requested the review has reached their quota limit.

Copilot

Copilot was unable to review this pull request because the user who requested the review has reached their quota limit.

Copilot

Pull request overview

Copilot reviewed 6 out of 6 changed files in this pull request and generated 1 comment.

shoumikhin · 2026-09-01T23:15:27Z

    if (is_capture_step) {
      // ----- CUDA graph CAPTURE -----
      ET_LOG(
          Info,
          "CUDA graph: beginning stream capture for '%s'",
          handle->method_name.c_str());

      cudaError_t cerr =
          cudaStreamBeginCapture(cuda_stream, cudaStreamCaptureModeRelaxed);
      ET_CHECK_OR_RETURN_ERROR(
          cerr == cudaSuccess,
          Internal,
          "cudaStreamBeginCapture failed: %s",
          cudaGetErrorString(cerr));
+      capture_guard.arm(cuda_stream, &handle->cuda_graph_state);
    }


Already handled at this head. The guard is armed before any of that work now, not after the capture begins: it is armed a few lines above the first cudaMalloc for the static inputs, so a failure while seeding them frees what was pinned and disables graphs for the method rather than leaking. Ending the capture is a second, separate guard, declared so it runs before the tensor cleanup, because freeing a buffer on a still-capturing stream fails and leaks the block.

Copilot

Pull request overview

Copilot reviewed 6 out of 6 changed files in this pull request and generated 2 comments.

shoumikhin · 2026-09-01T23:15:29Z

      } else if (std::strcmp(option.key, kUseSharedCudaStream) == 0) {
-        if (auto* val = std::get_if<bool>(&option.value)) {
-          if (*val) {
-            create_shared_cuda_stream();
-          }
-        } else {
-          ET_LOG(Error, "Option %s must be a boolean.", kUseSharedCudaStream);
-          return Error::InvalidArgument;
-        }
+        ET_LOG(
+            Info,
+            "Runtime backend option '%s' is DEPRECATED and has no effect. Methods "
+            "now run on the calling thread's stream, which orders methods called "


This is still validated. A non-boolean value falls to the else branch and returns InvalidArgument with a message saying the option must be a boolean. What changed is that a true value is now refused with NotSupported rather than honoured, because the option was the only thing ordering methods across threads and silently dropping it would give such a caller wrong results.

shoumikhin · 2026-09-01T23:15:31Z

+#define cudaStreamPerThread hipStreamPerThread
+


Fair, added. It is wrapped in an ifndef now, so a HIP header that defines the CUDA spelling itself will not warn.

Copilot

Copilot was unable to review this pull request because the user who requested the review has reached their quota limit.

Gasoonjia

It is really good. Thanks for the udpate.
I think my only concern is after this PR we need to put the graphs running parallel in GPU into different cpu thread but it should be ok.

Copilot

Copilot was unable to review this pull request because the user who requested the review has reached their quota limit.

shoumikhin · 2026-09-02T00:52:47Z

On the concern in your approval, you are right and it is worth writing down.

Two delegates on one thread now share that thread's stream, so they run one after another. Anything meant to overlap on the device has to come from separate host threads, one stream each, which is also what makes the ordering safe: within a thread the graph order is the device order.

There is an escape hatch for a caller who wants overlap without extra threads. A caller installed stream takes precedence over the per-thread one, so a caller can hand each call its own stream and drive concurrency itself. That path is unchanged by this change and stays the way to express deliberate overlap.

So nothing here blocks parallel execution, it moves where the decision is made: previously a stream per handle produced overlap nobody asked for and no ordering, now overlap is something a caller opts into by giving a stream or a thread.

Each delegate handle created its own CUDA stream. Delegates in one program run one after another, so a stream apiece bought no ordering between them: a delegate could read an input while the delegate that produces it still had work queued on a different stream. A program that is entirely one backend does not notice while it stays on one thread, because the stream is then the same throughout. A program split across the CUDA and TensorRT backends does. The TensorRT delegate runs on the per-thread stream when no caller stream is installed, so the two backends sat on different streams with nothing between them. Such a program returned a different answer on almost every call: executing one loaded program forty times produced nine distinct output sums, one of them correct. Every handle now takes the per-thread stream, which makes the ordering the graph expresses the ordering the device sees, in either direction and between two CUDA delegates, with no event handshake needed because both backends end up on the same stream. A caller stream still takes precedence per execute. The guarantee is per thread, and that is the one thing to get right about this change. cudaStreamPerThread resolves to a different stream on each host thread, so delegates called from one thread are ordered and delegates called from different threads are not. That is a real loss for one case. use_shared_cuda_stream created a single stream held by the backend and handed it to every handle on every thread, so a caller running the encoder on one thread and the decoder on another did get ordering from it. The option is now accepted and ignored, and its log line says plainly that ordering is per thread and that such a caller has to order the calls itself. Both in-repo callers drive their methods from one thread and no longer set it. Two smaller consequences of one stream rather than many. Two independent programs submitted from one thread with no synchronization between them now run one after the other rather than overlapping, measured at about 2x on two equal single-block kernels; a caller stream per program restores the overlap. And a handle no longer owns a stream that destroy could free while the thread-local table still pointed at it, which was a real dangling-stream path before. Because one stream is now shared, a capture that an error abandons would otherwise leave that stream capturing for every later delegate on the thread. A scope guard ends the capture and also frees the buffers the attempt pinned and resets the method to warmup, so a retry captures from a clean state. Without that second half the retry appends a second set of static buffers, replay reads one set while the input copies target the other, and every execute returns the values captured at that moment with nothing reporting an error, which is worse than the loud failure it replaced. Test plan: backends/cuda/tests/test_coalesced_determinism.py exports a program split across both backends, checks the saved program contains both backends, runs it a hundred times on one loaded program, and requires every result to equal the first exactly. It passes with this change and fails without it at the second run. Each step of the model carries its own buffer: one constant read by every island collapses to the same placeholder name several times in the flattened graph, which is a separate export defect and would fail this test before it reached the run loop. Verified separately on an H100, since no test covers it: a capture step that fails after capture began leaves the stream clean and the method's buffers freed, and three replays after the retry return the current answer. With the guard restoring only the stream, all three returned the capture-time value instead. Measured on Linux aarch64, sm_110, on a program of twenty-five delegates: before, thirty-nine of forty runs were wrong; after, seven thousand six hundred consecutive runs were correct across several processes. Programs of a single delegate on either backend were already correct and stayed so over five hundred runs each. Median latency for that split program went from about 850 to about 940 microseconds, which is the ordering that was previously skipped. A single delegate program is unchanged. The three existing C++ test binaries for this backend pass. The Python suite for this backend has eighteen failures on this machine both with and without this change, so it introduces no regressions; those are an unrelated gap in matmul and convolution lowering on this architecture. cudaStreamPerThread has no alias in the HIP compatibility header, so this adds one, or the ROCm build of this backend would not compile. That alias was checked by hand against a stub of the ROCm definitions and not by a ROCm build: the ROCm and Windows jobs are skipped for pull requests from a fork, so nothing in CI compiled it here. Not covered: the test above needs the TensorRT delegate, which nothing in the repo installs, so it is collected and skipped in CI, and that skip is silent because the job clears pytest's reporting flags. Installing the delegate there would not fix it either, because that job runs on x86 where this reordering does not reproduce; real coverage needs a job on an architecture where it does. Also not covered: Windows, and the CUDA graph capture and replay paths beyond confirming that a program using them runs and agrees with itself over forty runs.

Review found several problems with the previous revision. The deprecated option was accepted and ignored. That option was the only thing ordering methods driven from different threads, so a caller who set it for that reason kept building, kept running, and got unordered device work and wrong numbers. The one signal was an Info log, and a default Release configure compiles ET_LOG out entirely, so in practice there was no signal at all. It now returns NotSupported when set to true, matching the shape the graph option already uses for ROCm, and InvalidArgument for a non-boolean, which the previous revision had stopped checking. The long explanation moved into a comment: formatted into the runtime's 256 byte log buffer that message came to 338 characters and was cut mid-word, losing exactly the sentence telling the caller what to do. The capture teardown ran in the wrong order. The tensor cleanup guard was declared after the capture guard, so it was destroyed first and issued its frees into a stream that was still capturing. Measured on an H100: a free on a capturing stream returns invalid argument and the block is not returned, 64 MiB in the probe. Ending the capture is now its own guard, declared after the cleanup so it runs before it. With that order the unwind ends the capture, deletes the tensors, then frees, and nothing leaks. Three more, all on paths an error takes: The guard was disarmed before the first graph launch, the output copies and the final synchronize, so a failure in any of them left the method in the replay phase with the static output pointers recorded while the cleanup guard deleted those same buffers. The disarm moved past the last failure point, and the AOTI-owned outputs are released from the cleanup list before the copies rather than after. A static input buffer was tracked only after its seeding copy, so a failed copy leaked it. It is tracked as soon as the allocation succeeds. On capture failure the guard rewound to a full warmup, so a capture that can never succeed returned an error on one call in four for the life of the process. It now disables graphs for that method, which fails once and then runs eagerly. Also: the guard destroyed the graph before its exec, the opposite order from ~CudaGraphState, so the three teardown copies had already drifted; the handle's stream field defaulted to null, which is the legacy default stream and a genuinely different stream from the per-thread one the comment claimed, and it is now initialised to that sentinel so the invariant holds by construction; a comment claimed cudaGetLastError only clears this call's error, which it does not; a nine line comment in the ASR runner survived the code it described and explained export flags that file does not touch; the CMake block defining CUDA_AVAILABLE for the ASR runner lost its last consumer and the top level already defines it; and the HIP alias moved up with the other value aliases with a note on why it has to stay a macro. Test plan: The determinism test now compares the first result against eager before the run to run loop. Comparing a program only against itself passes a delegate that reads a stale buffer the same way every call, which is the same class of bug this change fixes. Verified on an H100, none of it covered by CI: a free into a still-capturing stream returns invalid argument and leaks; with the guards in this order the unwind ends the capture first and the free succeeds one executable graph replayed from two threads on two streams returns no error and serializes, so the refusal to mix a caller stream with graphs is about the shared static buffers and not a CUDA rule; the comment says that now the deprecation message formats to 338 characters into a 256 byte buffer hipStreamPerThread is ((hipStream_t)2), a cast to a pointer type, so it cannot be an inline constexpr like its neighbours Both guard classes were extracted and compiled standalone with nvcc, and the handle's new default was checked to equal cudaStreamPerThread (0x2) where the old one was null. clang-format is clean on all four C++ files; black and usort are clean on the test. The full backend still has no CI job that builds it from a fork, so this is not a build of the backend itself.

The HIP compatibility section defines cudaStreamPerThread unconditionally, so a HIP header that defines the CUDA spelling itself would warn. Guarded.

Copilot

Copilot was unable to review this pull request because the user who requested the review has reached their quota limit.

meta-cla Bot added the CLA Signed This label is managed by the Facebook bot. Authors need to sign the CLA before a PR can be reviewed. label Aug 30, 2026

shoumikhin force-pushed the fix/cuda-delegate-per-thread-stream branch 3 times, most recently from 7a149ef to 315df2c Compare August 30, 2026 18:03

shoumikhin added the release notes: runtime Changes related to the core runtime which loads the program methods, initializes delegates, and runs label Aug 30, 2026

shoumikhin force-pushed the fix/cuda-delegate-per-thread-stream branch from 315df2c to 95b8741 Compare August 30, 2026 19:27

shoumikhin marked this pull request as ready for review August 30, 2026 21:08

shoumikhin requested review from kirklandsign and larryliu0820 as code owners August 30, 2026 21:08

Copilot AI lite review requested due to automatic review settings August 30, 2026 21:08

Copilot AI reviewed Aug 30, 2026

shoumikhin force-pushed the fix/cuda-delegate-per-thread-stream branch from 95b8741 to 701747b Compare August 31, 2026 04:30

Copilot AI review requested due to automatic review settings August 31, 2026 04:30

Copilot AI reviewed Aug 31, 2026

shoumikhin force-pushed the fix/cuda-delegate-per-thread-stream branch from 701747b to 7d2ce28 Compare August 31, 2026 06:01

Copilot AI review requested due to automatic review settings August 31, 2026 06:01

Copilot AI reviewed Aug 31, 2026

shoumikhin force-pushed the fix/cuda-delegate-per-thread-stream branch from 7d2ce28 to 80676cc Compare August 31, 2026 15:20

Copilot AI review requested due to automatic review settings August 31, 2026 15:20

Copilot AI reviewed Aug 31, 2026

shoumikhin force-pushed the fix/cuda-delegate-per-thread-stream branch from 80676cc to 8a1c058 Compare August 31, 2026 22:09

Copilot AI review requested due to automatic review settings August 31, 2026 22:09

Copilot AI reviewed Aug 31, 2026

shoumikhin force-pushed the fix/cuda-delegate-per-thread-stream branch from 8a1c058 to d8f8bcb Compare August 31, 2026 22:16

Copilot AI review requested due to automatic review settings August 31, 2026 22:16

Copilot AI reviewed Aug 31, 2026

shoumikhin force-pushed the fix/cuda-delegate-per-thread-stream branch from d8f8bcb to 0a3f76f Compare September 1, 2026 17:00

Copilot AI review requested due to automatic review settings September 1, 2026 17:00

Copilot started reviewing on behalf of shoumikhin September 1, 2026 17:00 View session

Copilot AI reviewed Sep 1, 2026

View reviewed changes

Copilot AI review requested due to automatic review settings September 1, 2026 17:14

shoumikhin force-pushed the fix/cuda-delegate-per-thread-stream branch from 0a3f76f to e5e7564 Compare September 1, 2026 17:14

Copilot started reviewing on behalf of shoumikhin September 1, 2026 17:15 View session

Copilot AI reviewed Sep 1, 2026

View reviewed changes

Copilot AI review requested due to automatic review settings September 1, 2026 22:21

Copilot AI reviewed Sep 1, 2026

Gasoonjia approved these changes Sep 1, 2026

View reviewed changes

Copilot AI review requested due to automatic review settings September 1, 2026 23:15

Copilot AI reviewed Sep 1, 2026

Anthony Shoumikhin and others added 3 commits September 1, 2026 17:55

Guard the per-thread stream macro against a redefinition

374821d

The HIP compatibility section defines cudaStreamPerThread unconditionally, so a HIP header that defines the CUDA spelling itself would warn. Guarded.

Copilot AI review requested due to automatic review settings September 2, 2026 00:55

shoumikhin force-pushed the fix/cuda-delegate-per-thread-stream branch from f823796 to 374821d Compare September 2, 2026 00:55

Copilot AI reviewed Sep 2, 2026

shoumikhin merged commit 0334561 into pytorch:main Sep 2, 2026
238 of 242 checks passed

Uh oh!

Conversation

shoumikhin commented Aug 30, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

pytorch-bot Bot commented Aug 30, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

🔗 Helpful Links

🧪 See artifacts and rendered test results at hud.pytorch.org/pr/pytorch/executorch/22318

❌ 3 New Failures, 2 Unrelated Failures

Uh oh!

linux-foundation-easycla Bot commented Aug 30, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Uh oh!

shoumikhin Sep 1, 2026

Choose a reason for hiding this comment

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Uh oh!

shoumikhin Sep 1, 2026

Choose a reason for hiding this comment

Uh oh!

shoumikhin Sep 1, 2026

Choose a reason for hiding this comment

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Uh oh!

Gasoonjia left a comment

Choose a reason for hiding this comment

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Uh oh!

shoumikhin commented Sep 2, 2026

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

shoumikhin commented Aug 30, 2026 •

edited

Loading

pytorch-bot Bot commented Aug 30, 2026 •

edited

Loading

linux-foundation-easycla Bot commented Aug 30, 2026 •

edited

Loading