Repository navigation
perf(rccl): skip the temporary stream when capturing single-stream launches - #14118
Conversation
…unches During graph capture, ncclLaunchFinish fast-forwards deviceStream to the launch stream with ncclStreamAdvanceToEvent, which creates a stream, waits on finishedEvent there, copies its capture dependencies onto deviceStream and destroys it again. On ROCm 7.0.2 the create and destroy cost about 1.5 ms per captured launch: 16 captured all_reduces spend 3.9 ms in hipStreamCreate and 19.7 ms in hipStreamDestroy out of 23.9 ms total. When the launch has one user stream, finishedEvent has just been recorded on it and nothing is captured on it before the fast-forward, so its current capture dependencies are exactly what a wait on finishedEvent would install. For that case, add ncclStreamAdvanceToStream, which reads them with hipStreamGetCaptureInfo and sets them on deviceStream directly. Uncaptured launches and launches over several user streams still use ncclStreamAdvanceToEvent. Its dependency-copy code, including the CUDA 12.3 edge-data and CUDA 13 branches, moves into a static helper that both functions share. The change applies to the default NCCL_GRAPH_STREAM_ORDERING=1; with ordering 0, captured launches run on the launch stream and never fast-forward deviceStream. 8 x MI350X, 2 ranks, in-place uint32 all_reduce, capture time per collective (old -> new): 1 collective 1.44-1.62 ms -> 0.012-0.033 ms; 16 collectives 1.42-1.62 ms -> 0.009 ms; 64 collectives 1.62-1.70 ms -> 0.008 ms. The first capture on a communicator still costs about 1 ms. With the change, captures make no hipStreamCreate or hipStreamDestroy calls; before, they made two of each per captured launch. Node and edge counts and replay results are unchanged, with and without NCCL_GRAPH_MIXING_SUPPORT=1. Add GraphCaptureStreamOrderingTests (11 tests, 29 process-isolated configurations on 2 GPUs). Each test captures a sequence, checks the graph's node types and which nodes can reach which, bounds how many dependencies any node has, then replays the graph and checks the data. With graph mixing, it also checks that every collective reaches a serialEvent record, which ncclLaunchFinish adds when it releases deviceStream. They cover: - one collective under all three capture modes, and chains of 16 and 128 - user kernels between collectives under each capture mode, with graph mixing and implicit launch order on and off, and over the socket transport, where plans add host-callback nodes - fork/join user work feeding a collective, and grouped collectives plus send/recv on one stream - uncaptured launches before, between and after graph launches, including a two-stream group, with implicit launch order, with graph mixing and with NCCL_GRAPH_STREAM_ORDERING=0 - two graphs captured at once from the same communicators, recapture on the same and on a different stream, a single-rank communicator, and zero-count collectives The suite passes with and without this change. Replacing the fast-forward with a plain hipStreamWaitEvent, or adding to the dependency set instead of replacing it, fails 6 configurations through the dependency bound. Skipping the fast-forward fails the 3 graph-mixing configurations without implicit launch order through the serialEvent check. Not covered: - Captures with several user streams: on ROCm 7.0.2 they crash in hipStreamEndCapture even without this change. - NCCL_GRAPH_STREAM_ORDERING=0 when user work precedes the first captured collective: on ROCm 7.0.2, hipStreamWaitEvent with hipEventWaitExternal segfaults once the capturing stream has any node, so rcclGraphOriginWaitSerialEvent crashes there even without this change. With NCCL_LAUNCH_ORDER_IMPLICIT=1, captured kernels depend on every earlier kernel, because ncclLaunchFinish makes launchOrder plain-wait on each launch. That is not changed here, so those configurations check ordering and data only. Co-authored-by: Cursor <cursoragent@cursor.com>
…ROCm 7.0.2 With NCCL_GRAPH_STREAM_ORDERING=0, ncclLaunchPrepare makes the capturing launch stream wait on deviceStream.serialEvent with hipEventWaitExternal (rcclGraphOriginWaitSerialEvent). On ROCm 7.0.2 that call segfaults in hip::Graph::AddExternalEventWaitNode once the stream has captured any node, so capturing a collective after a memset or a user kernel crashes. A capture that starts with the collective works, which is why EagerAroundCapture/StreamOrdering0 passes. A standalone HIP program reproduces the crash with one captured memset followed by an external wait, so the bug is in HIP and RCCL only exposes it. Develop covers graphStreamOrdering=0 only with host fakes and MPI tests, so CI does not hit this. Add GraphCaptureStreamOrdering.StreamOrdering0UserWorkBeforeCollective with three configurations: one memset and one all_reduce, and the interleaved user-work capture with implicit launch order off and on. It skips unless RCCL_TEST_KNOWN_FAILURES=1. A DISABLED_ prefix would not work: the process-isolated runner re-execs children with only --gtest_filter, so they would run nothing and report a pass. With the variable set, all three configurations die with signal 11. Co-authored-by: Cursor <cursoragent@cursor.com>
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
There was a problem hiding this comment.
ARGUS
Findings
4 more finding(s), not posted inline
projects/rccl/src/enqueue/enqueue.cc:2903(Low) ncclStreamAdvanceToStream has no fail-loud stub in test/host/fakes/strongstream_stubs.cc, unlike ncclStrongStreamAcquiredWorkStream (strongstream_stubs.cc:57-60); add FailLoudUnfaked("strongstream_stubs", "ncclStreamAdvanceToStream") beside it.projects/rccl/src/misc/strongstream.cc:420(Low) ncclStreamAdvanceToStream only forwards to static setCaptureDependenciesFrom; define the public function directly and call it from ncclStreamAdvanceToEvent.projects/rccl/test/GraphCaptureStreamOrderingTests.cpp:394(Nit) Reflow the comment at lines 390-394; line 391 is about 157 columns.projects/rccl/test/GraphCaptureStreamOrderingTests.cpp:22(Nit) Add <chrono> for std::chrono::seconds at line 500, as DoneEventOrderingTests.cpp:14 does.
This PR changes 1,017 lines across 5 files. A change this size is hard to review
well, for a person and for this bot alike: reviewers lose the thread, defects hide in the
volume, and the feedback arrives too late to be cheap to act on. Splitting it into a
stacked pull request
gives each piece a reviewable diff and keeps the history. This review covers the whole change
regardless.
Bottom line: Worth settling before approval: GraphCaptureStreamOrdering.* runs in no CI config and ncclStreamAdvanceToStream has no microtest, so the new path is untested in CI; also the UB shift for Chain128 at GraphCaptureStreamOrderingTests.cpp:443.
Comment @argus review to re-run, or @argus re-review after pushing.
- Run GraphCaptureStreamOrdering.* in CI: add each test to ci-precheckin.json next to DoneEventOrdering, and add the suite to the standard/comprehensive/full patterns in test_categories.yaml. - Add host microtests for ncclStreamAdvanceToStream: it reads the capture info of `source`, SETs the dependencies of `s`, creates and destroys no stream, and propagates HIP failures. - Add a fail-loud ncclStreamAdvanceToStream stub to strongstream_stubs.cc for the host binaries that link enqueue.cc. - Make ncclStreamAdvanceToStream the dependency-copy function itself instead of a wrapper around a static helper; ncclStreamAdvanceToEvent calls it with its scratch stream. - Compute the expected Chain128 value with wrapping arithmetic; the old `3ull << (n - 1)` shifted a 64-bit value by 127. - Label host-callback nodes and require some in InterleavedUserWorkNetTransport, so it checks that the socket plans really carry proxy ops. - Shorten the ncclLaunchFinish comment, reflow a long test comment and include <chrono>. Co-authored-by: Cursor <cursoragent@cursor.com>
There was a problem hiding this comment.
ARGUS
This PR changes 1,062 lines across 9 files. A change this size is hard to review
well, for a person and for this bot alike: reviewers lose the thread, defects hide in the
volume, and the feedback arrives too late to be cheap to act on. Splitting it into a
stacked pull request
gives each piece a reviewable diff and keeps the history. This review covers the whole change
regardless.
Bottom line: LGTM.
Comment @argus review to re-run, or @argus re-review after pushing.
|
The description still mentions a static |
…am capture PR
Source:
- rcclAddonCaptureOrderEnd records finishedEvent on the launch stream right
before the fast-forward, so under capture it now copies the launch
stream's frontier with ncclStreamAdvanceToStream instead of
ncclStreamAdvanceToEvent. Small DDA and copy-engine collectives take this
path, so they no longer create and destroy a stream per captured launch.
- ncclStreamAdvanceToStream: initialize `nodes`. HIP fills it only for an
active capture, so on an invalidated one it reached
hipStreamUpdateCaptureDependencies unset (with count 0).
Host microtests:
- ncclStreamAdvanceToStream and ncclStrongStreamAcquiredWorkStream are
hookable in the strongstream fakes, with ASSERT_HOOK_MATCHES_PROD checks
and reset in ResetStrongStreamStubs.
- New EnqueueMicrotests drive ncclLaunchFinish while capturing: one user
stream must call AdvanceToStream(deviceStream, launchStream) and not
AdvanceToEvent; two user streams must call AdvanceToEvent. Reverting the
branch to AdvanceToEvent now fails the single-stream test.
- The addon capture tests expect AdvanceToStream in place of
record + AdvanceToEvent.
GraphCaptureStreamOrderingTests (9 tests, 21 configurations, was 11 / 33):
- expectBuffer uses RCCLTestHelpers::verifyBufferData; withHarness and
replayAndExpect replace the repeated Harness/initHarness/replay blocks.
- Collective counts use count('C') >= n and orders use collapse(), so a
protocol or copy-engine change that emits more than one kernel node per
collective does not fail them.
- Drop Chain128: Chain16 already catches add-vs-set via redundantEdges().
- InterleavedUserWork: capture mode does not reach the new branch, so run
graph mixing x implicit order under ThreadLocal plus one Global child
(5 children instead of 12). Each round also issues a zero-count
AllReduce, which builds no plan; this replaces CallsWithoutKernelPlan.
- EagerMultiStreamAroundCapture is now the TwoStreamEager option of
EagerAroundCapture.
- Drop StreamOrdering0UserWorkBeforeCollective and its env-var skip. With
NCCL_GRAPH_STREAM_ORDERING=0 the new branch is never reached; the HIP
segfault it documented is to be tracked in a ticket.
Not added: a collective launched on a stream forked inside the capture.
capturingGraph.origin is the user stream passed to the collective (set in
ncclPlannerSetCapturingGraph), so with one stream it is launchStream and
reading the frontier from it cannot be told apart. Launching on a forked
stream also crashes on ROCm 7.0.2 with or without this change: HIP's
hip::Stream::EndCapture recurses until the stack overflows when a stream
joins the capture through a non-origin stream that then waits on it,
which is what ncclStrongStreamAcquire + ncclLaunchPrepare do there.
Co-authored-by: Cursor <cursoragent@cursor.com>
There was a problem hiding this comment.
ARGUS
This PR changes 1,069 lines across 11 files. A change this size is hard to review
well, for a person and for this bot alike: reviewers lose the thread, defects hide in the
volume, and the feedback arrives too late to be cheap to act on. Splitting it into a
stacked pull request
gives each piece a reviewable diff and keeps the history. This review covers the whole change
regardless.
Bottom line: LGTM.
Comment @argus review to re-run, or @argus re-review after pushing.
This comment was marked as off-topic.
This comment was marked as off-topic.
This comment was marked as off-topic.
This comment was marked as off-topic.
|
@pvallem Thanks for the review, all comments are addressed in 3cd8361. I updated the description: the shared-code bullet now says |
Updated. The "Shared code" bullet now says |
There was a problem hiding this comment.
🟢 Approval recommended
No active review issues block approval of the capture-dependency optimization.
No open issues 🎉
What changed in this PR
This PR reduces RCCL’s HIP graph capture overhead by copying capture dependencies directly from the launch stream instead of creating and destroying a temporary stream for single-stream captured launches.
Changes:
- Adds a shared capture-dependency helper and uses it in single-stream launch completion and addon capture paths, preserving existing fallback behavior.
- Adds host microtests and GPU tests for graph ordering, dependency structure, replay results, and capture lifecycle scenarios.
- Registers the GPU suite in the test build, test categories, and CI precheckin configuration.
| File | Description |
|---|---|
projects/rccl/tools/scripts/test_runner/configs/ci-precheckin.json |
Adds nine graph-capture tests to precheckin runs. |
projects/rccl/test/test_categories.yaml |
Categorizes the graph-capture test suite. |
projects/rccl/test/host/strongstream-test.cc |
Tests dependency replacement and HIP error propagation. |
projects/rccl/test/host/fakes/strongstream_stubs.h |
Declares new strong-stream test hooks. |
projects/rccl/test/host/fakes/strongstream_stubs.cc |
Implements and resets the new hooks. |
projects/rccl/test/host/enqueue-test.cc |
Tests single-stream optimization, multi-stream fallback, and addon behavior. |
projects/rccl/test/GraphCaptureStreamOrderingTests.cpp |
Adds GPU topology and replay tests across capture scenarios. |
projects/rccl/test/CMakeLists.txt |
Includes the new GPU tests in the build. |
projects/rccl/src/misc/strongstream.cc |
Implements direct dependency copying and shares it with the event-based path. |
projects/rccl/src/include/strongstream.h |
Declares the capture-only helper and documents its preconditions. |
projects/rccl/src/enqueue/enqueue.cc |
Applies direct dependency copying to eligible captured launches. |
Review effort: Balanced
TheRock Submodule Bump ActivityNewest first
|
perf(rccl): skip the temporary stream when capturing single-stream launches
Motivation
Capturing an RCCL collective into a HIP graph costs about 1.5 ms per collective on ROCm 7.0.2, almost all of it outside RCCL's own work. During capture,
ncclLaunchFinishfast-forwardsdeviceStreamto the launch stream withncclStreamAdvanceToEvent. That function creates a temporary stream, waits onfinishedEventthere, copies the temporary stream's capture dependencies ontodeviceStreamand destroys it again. For 16 captured all_reduces, 3.9 ms goes tohipStreamCreateand 19.7 ms tohipStreamDestroy, out of 23.9 ms total.This PR removes the temporary stream for the common case: a captured launch on a single user stream, under the default
NCCL_GRAPH_STREAM_ORDERING=1. The same applies to the addon path that small DDA and copy-engine collectives take. In both places RCCL now reads the capture dependencies straight from the launch stream. Graph structure and replay results are unchanged. Uncaptured launches, launches over several user streams andNCCL_GRAPH_STREAM_ORDERING=0keep their current paths.Technical Details
Fast-forward without a temporary stream
finishedEventhas just been recorded on the launch stream, and nothing is captured on that stream before the fast-forward. The launch stream's current capture dependencies are therefore exactly what a wait onfinishedEventwould install on another stream.ncclStreamAdvanceToStream(s, source)(strongstream.h/.cc). It readssource's capture dependencies withhipStreamGetCaptureInfoand sets them onswithhipStreamUpdateCaptureDependencies(..., SET). Both streams must be capturing into the same graph. It is capture-only; uncaptured work still needs a real event wait.ncclStreamAdvanceToEventnow creates its temporary stream, waits on the event there, and callsncclStreamAdvanceToStreamwith it, so both paths use one copy of the dependency code. That includes the CUDA 12.3 lossy-query edge-data branch and the CUDA 13_v3/_v2branch.src/enqueue/enqueue.cc):ncclLaunchFinish: inside develop'sif (!useLaunchStream)block,capturing && planner->numStreams == 1now callsncclStreamAdvanceToStream(deviceStream, launchStream). Everything else still callsncclStreamAdvanceToEvent. WithNCCL_GRAPH_STREAM_ORDERING=0(useLaunchStream), captured kernels run on the launch stream and never fast-forwarddeviceStream, so that path is untouched.rcclAddonCaptureOrderEnd(DDA and copy-engine collectives under capture):ncclStreamAdvanceToStream(deviceStream, stream)replaces recordingscratchEventonstreamand then callingncclStreamAdvanceToEvent.Tests
Host microtests (
test/host/enqueue-test.cc,test/host/strongstream-test.cc):ncclStreamAdvanceToStreamandncclStrongStreamAcquiredWorkStreamare hookable in the strongstream fakes, likencclStreamAdvanceToEvent.LaunchFinish_CapturedSingleStream_*andLaunchFinish_CapturedTwoStreams_*drivencclLaunchFinishwhile capturing. One user stream must callAdvanceToStream(deviceStream, launchStream)and neverAdvanceToEvent; two user streams must callAdvanceToEvent.AdvanceToStreamin place of the event record andAdvanceToEvent.AdvanceToStream_*tests check thatncclStreamAdvanceToStreamreplaces the frontier with the source's and propagates HIP failures.GPU tests.
GraphCaptureStreamOrderingTests.cppadds 9 tests with 21 process-isolated configurations on 2 GPUs.Each configuration captures a sequence of RCCL and user work and then checks:
NCCL_GRAPH_MIXING_SUPPORT=1, that every collective reaches aserialEventrecord node, whichncclLaunchFinishadds when it releasesdeviceStream;The configurations cover:
NCCL_GRAPH_STREAM_ORDERING=0;Known issues, not fixed
These all reproduce on develop without this change.
NCCL_GRAPH_STREAM_ORDERING=0with user work before the first captured collective segfaults on ROCm 7.0.2.rcclGraphOriginWaitSerialEvent(andrcclAddonCaptureOrderBegin) callhipStreamWaitEvent(..., hipEventWaitExternal), and HIP crashes inhip::Graph::AddExternalEventWaitNodeonce the capturing stream has any node. A short HIP program reproduces it without RCCL. A HIP ticket is to be filed.hipStreamEndCaptureon ROCm 7.0.2.ncclStrongStreamAcquirejoinsdeviceStreaminto the capture through the launch stream, and the launch stream then waits ondeviceStream. When the launch stream is not the capture origin, HIP'ship::Stream::EndCapturerecurses until the stack overflows. A HIP-only program with the same three-stream pattern reproduces it. A HIP ticket is to be filed.hipStreamEndCaptureon ROCm 7.0.2, in every capture mode. Captured multi-stream launches are covered by the host microtests only.NCCL_LAUNCH_ORDER_IMPLICIT=1, every captured kernel depends on all earlier ones, becauselaunchOrderplain-waits on each launch. Those configurations check ordering and data only.Issue Tracking
JIRA ID: ROCM-33215 (Confluence: MLSE "Issue 3: hipStreamCreate Cost in the RCCL Collective Capture Path")
Test Plan
One node, 8 × MI350X (gfx950), ROCm 7.0.2.2. RCCL and
rccl-UnitTestswere built (Release, gfx950) from this branch on develope790e3fef07.GraphCaptureStreamOrdering.*, all 21 configurations, against:hipStreamWaitEvent;rccl-UnitTestsMicro*binary; for the enqueue microtests, also a mutant that reverts thencclLaunchFinishbranch toncclStreamAdvanceToEvent.3cd83617f0a.Test Result
GPU tests. The suite passes with and without the fix, and each mutant fails at least one configuration:
serialEventcheckhipStreamWaitEventinsteadHost microtests. All pass except
NcclDeviceRequirementsMicrotest.GinInboxA2ACreateRequirement_KnownBug_SingleRankTeamDividesByZero, which this PR does not touch. Reverting thencclLaunchFinishbranch failsLaunchFinish_CapturedSingleStream_AdvancesDeviceStreamToTheLaunchStream.2-node rccl-tests. All 66 runs passed with 0 wrong values.
Capture time per collective, before → after:
Before, each captured launch made 2
hipStreamCreateand 2hipStreamDestroycalls; now it makes none. The first capture on a communicator still costs about 1 ms of one-time setup.