Repository navigation
feat(hip): enable stream capture for hipStreamWaitValue/WriteValue/Ba… - #8911
Merged
Merged
Conversation
❌ PR Check — Action Required
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🚫 Please fix the failed policies before requesting reviews. The following policy checks failed:
The |
saleelk
reviewed
Jul 21, 2026
anugodavar
force-pushed
the
users/agodavar/batchmemops
branch
2 times, most recently
from
July 21, 2026 12:03
6fccc5e to
f9a0ddc
Compare
anugodavar
marked this pull request as ready for review
July 21, 2026 12:06
Contributor
There was a problem hiding this comment.
Pull request overview
This PR extends HIP graph stream-capture support to include hipStreamWaitValue*, hipStreamWriteValue*, and hipStreamBatchMemOp, ensuring these APIs record BatchMemOp graph nodes during capture rather than executing immediately. It also fixes correctness and replay-safety issues in BatchMemOp graph nodes and the underlying blit-kernel execution ordering.
Changes:
- Add stream-capture handlers so wait/write value ops and batch mem ops capture as BatchMemOp graph nodes, matching CUDA behavior.
- Fix
hipGraphBatchMemOpNodeto deep-copyparamArrayto avoid dangling pointers on graph replay/clone. - Enforce sequential BatchMemOp execution order (CUDA spec) by launching a single work-item and looping in the kernel, plus add/extend unit tests.
Reviewed changes
Copilot reviewed 13 out of 13 changed files in this pull request and generated 2 comments.
Show a summary per file
| File | Description |
|---|---|
| projects/hip-tests/catch/unit/stream/hipStreamValueCapture.cc | New unit tests validating stream-capture behavior for wait/write value and batch mem ops. |
| projects/hip-tests/catch/unit/stream/hipStreamBatchMemOp.cc | Add sequencing stress tests (single- and multi-GPU) to validate sequential batch execution. |
| projects/hip-tests/catch/unit/stream/CMakeLists.txt | Include the new stream-capture test in the stream unit test build. |
| projects/hip-tests/catch/unit/graph/hipGraphBatchMemOpNodeCapture.cc | New unit test for BatchMemOp graph node instantiate/launch behavior. |
| projects/hip-tests/catch/unit/graph/CMakeLists.txt | Include the new graph BatchMemOp node capture test in the graph unit test build. |
| projects/hip-tests/catch/config/configs/unit/stream.yaml | Register/disable/tag new stream tests in the unit test config. |
| projects/hip-tests/catch/config/configs/unit/graph.yaml | Register/disable/tag new graph test in the unit test config. |
| projects/clr/rocclr/device/rocm/rocblit.cpp | Enforce sequential BatchMemOp execution (globalWorkSize=1) and use stable kernarg allocation during graph packet capture. |
| projects/clr/rocclr/device/blitcl.cpp | Inline a sequential loop implementation of the BatchMemOp kernel until upstream amdblit.cl is fixed. |
| projects/clr/hipamd/src/hip_stream_ops.cpp | Route wait/write value and batch mem ops through STREAM_CAPTURE so they become graph nodes under capture. |
| projects/clr/hipamd/src/hip_graph.cpp | Implement capture handlers that create BatchMemOp nodes for captured stream wait/write and batch ops. |
| projects/clr/hipamd/src/hip_graph_internal.hpp | Deep-copy hipBatchMemOpNodeParams::paramArray to prevent dangling pointers across replay/clone. |
| projects/clr/hipamd/src/hip_graph_capture.hpp | Declare new capture handler entry points for wait/write value and batch mem ops. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
saleelk
approved these changes
Jul 21, 2026
eble-amd
reviewed
Jul 22, 2026
…tchMemOp - Add STREAM_CAPTURE hooks to hipStreamWaitValue32/64, hipStreamWriteValue32/64, and hipStreamBatchMemOp so they record BatchMemOp graph nodes during capture instead of executing immediately. - Declare and implement capture handlers in hip_graph_capture.hpp/hip_graph.cpp. - Fix hipGraphBatchMemOpNode to deep-copy paramArray to prevent dangling pointer on graph replay. - Fix batchMemOps() globalWorkSize from count to 1 so ops execute sequentially in array order per CUDA spec. - Inline sequential for-loop in __amd_rocclr_batchMemOp in blitcl.cpp to be self-contained until the fix lands in upstream llvm-project amdblit.cl. Co-Authored-By: Claude <noreply@anthropic.com>
…BatchMemOp Add Unit_hipStreamWaitWriteValue32_Capture, Unit_hipStreamWaitWriteValue64_Capture, and Unit_hipStreamBatchMemOp_Capture to verify that hipStreamWaitValue32/64, hipStreamWriteValue32/64, and hipStreamBatchMemOp are correctly captured as BatchMemOp graph nodes during stream capture and produce correct results after hipGraphLaunch. Co-Authored-By: Claude <noreply@anthropic.com>
hipMallocSignalMemory requires HSA signal memory (ROCr/Linux only). Co-Authored-By: Claude <noreply@anthropic.com>
- Add comment explaining mask is dropped on capture path: hipStreamBatchMemOpParams has no mask field; CUDA ignores mask universally (mask is AMD-only). - Clarify blitcl.cpp enum/struct definitions with source reference comment. Note: hipStreamWaitValue flags (0=GTE,1=EQ,2=AND,3=NOR) match compareOp values in __amd_streamOpsWait directly — no conversion needed. - Add Unit_hipStreamBatchMemOp_SequentialOrdering test: verifies ops execute in array order without graph (Write(1), Wait(==1), Write(0) — the race condition that motivated the globalWorkSize=1 fix). - Fix streamWaitValueSupported() in hipStreamValueCapture.cc (was missing helper). Co-Authored-By: Claude <noreply@anthropic.com>
… PR #6800) - Add GraphCaptureEnabled() override to hipGraphBatchMemOpNode so the blit kernel packet is pre-built at hipGraphInstantiate and dispatched via the flat AQL packet path at launch, matching memset/memcpy node behavior. - In batchMemOps(), allocate kernarg buffer from the graph's stable GraphKernelArgManager pool during AQL capture so the address baked into the captured AQL packet remains valid on re-launch (stream buffer is recycled between launches). - Add Unit_hipGraphBatchMemOpNode_AQLCapture test. Co-Authored-By: Claude <noreply@anthropic.com>
- Rename Unit_hipGraphBatchMemOpNode_AQLCapture to Unit_hipGraphBatchMemOpNode_WriteAndWait to accurately reflect what the test verifies (functional correctness, not AQL internals). - Add Unit_hipStreamBatchMemOp_SequentialOrdering: verifies Write(1)/Wait(==1)/Write(0) batch executes in array order on a single GPU. Fails without globalWorkSize=1 fix. - Add Unit_hipStreamBatchMemOp_SequentialOrdering_MultiGPU: same test across all GPUs, mirrors the original RCCL multi-GPU reproducer. Fails without the fix. - Add streamWaitValueSupported() helper to hipStreamBatchMemOp.cc. - Disable all new signal-memory tests on amd_windows, amd_wsl, nvidia. Co-Authored-By: Claude <noreply@anthropic.com>
hipExtMallocWithFlags with hipMallocSignalMemory is AMD-only and fails to compile on NVIDIA. Wrap the AMD-specific alloc/free calls with #if HT_AMD following the pattern used in hipStreamValue.cc. The runtime HIP_SKIP_TEST guard handles NVIDIA at execution time. Co-Authored-By: Claude <noreply@anthropic.com>
anugodavar
force-pushed
the
users/agodavar/batchmemops
branch
from
July 22, 2026 13:47
e5ee427 to
fb042fb
Compare
anugodavar
force-pushed
the
users/agodavar/batchmemops
branch
from
July 22, 2026 13:49
fb042fb to
5c9d6be
Compare
Contributor
Author
|
unit tests are added |
Comment on lines
+3792
to
+3794
| hipError_t capturehipStreamBatchMemOp(hipStream_t& stream, unsigned int& count, | ||
| hipStreamBatchMemOpParams*& paramArray, | ||
| unsigned int& flags) { |
Contributor
There was a problem hiding this comment.
Why do we need to accept everything by reference here? hipStream_t is an opaque pointer, plain integers take less space is passed directly, etc.
| hip::Stream* s = reinterpret_cast<hip::Stream*>(stream); | ||
|
|
||
| hipBatchMemOpNodeParams nodeParams{}; | ||
| nodeParams.ctx = reinterpret_cast<hipCtx_t>(hip::getCurrentDevice()); |
Contributor
There was a problem hiding this comment.
Is it a valid operation? hip::getCurrentDevice() returns Device * and I do not see why it would be interchangeable with context. Device class does have asContext() method though
1 task
AbandiGa
added a commit
that referenced
this pull request
Aug 31, 2026
## Motivation RCCL's CE data-movement collectives (AllGather, AlltoAll, Scatter, Gather) already run under HIP graph capture on develop, but the UC cross-rank barrier was not safe to replay: peer flag-writes, waits, and capture resets could share one hipStreamBatchMemOp. On HIP that batch does not guarantee wait-before-reset ordering, so a fused reset can clobber a peer flag mid-barrier and hang on graph replay. This PR makes the UC CE barrier graph-replay-safe on HIP, while keeping the fast fused write+wait path for non-capture launches (full NCCL-style always-split peer writes regress small-message latency) ## Technical Details Changes are in ce_coll.cc / ce_coll.h, plus whitebox test updates. No enqueue.cc changes; develop already routes data-movement CE under capture. 1. **Hybrid UC peer writes (ncclPrepUCSync)** - Non-capture: fuse peer WRITE_VALUE + WAIT_VALUE into one batchParams submit (same latency shape as develop). - Capture: issue peer writes as a separate hipStreamBatchMemOp first, then put only waits in batchParams, so the captured stream orders write → wait → reset. Peer values use immediate WRITE_VALUE with GRAPH_SYNC_VALUE while capturing (and the running seq otherwise). We do not use NCCL's D2D memcpy/ceSeqNumDev peer-write path for UC, to avoid a permanent non-graph latency tax. 3. **Even CE intra-batch sync count (ncclCeLaunchBatchOps)** Ported from NCCL: an odd sync count leaves useCompletePtr flipped and desyncs the next ready/complete barrier on replay; append one extra sync when the count would be odd. 4. **Graph-capture window creation** Port NCCL's fix ([22bf876](NVIDIA/nccl@22bf876)) so window creation stays graph-safe during capture. 5. **Tests** Update CeInternalMPITests.cpp for the new ncclPrepUCSync(..., stream) signature (hipStream_t). Non-capture whitebox expectations remain 2*(nRanks-1) ops in the batch. The MC path (`ncclPrepMCSync`) already issues its write separately and is unchanged. ## Issue Tracking JIRA ID: AICOMRCCL-1518 ## Test Plan Validation was performed with and without `NCCL_GRAPH_MIXING_SUPPORT=1`, and with `-G 20`, `-R 2`, and `-x 2`, successfully exercising graph capture and replay for CE collectives. To validate the CE graph path, the current graph guard that prevents RCCL from taking the CE path when graph capture is enabled was temporarily removed. This change was used for validation purposes only and is not included in this PR. The existing graph guard remains in `develop` and is not modified by this PR. CE Graph support will be re-enabled in a follow-up PR once the required HIP stream capture fix #8911 is available in an official ROCm release. At that point, the guard will be converted to ROCm version-based gating, with an environment-variable override if needed. Because the graph guard currently prevents this code path from being exercised in `develop`, no automated **test case is added as part of this PR.** ## Test Result Tested in MI355 with CLR build that include this PR #8911 and in MI450 with latest TheRock nightly build that include this fix. <!-- Pending. --> ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. --------- Co-authored-by: AbandiGa <abandiga@gmail.com>
corey-derochie-amd
pushed a commit
that referenced
this pull request
Oct 14, 2026
## Motivation RCCL's CE data-movement collectives (AllGather, AlltoAll, Scatter, Gather) already run under HIP graph capture on develop, but the UC cross-rank barrier was not safe to replay: peer flag-writes, waits, and capture resets could share one hipStreamBatchMemOp. On HIP that batch does not guarantee wait-before-reset ordering, so a fused reset can clobber a peer flag mid-barrier and hang on graph replay. This PR makes the UC CE barrier graph-replay-safe on HIP, while keeping the fast fused write+wait path for non-capture launches (full NCCL-style always-split peer writes regress small-message latency) ## Technical Details Changes are in ce_coll.cc / ce_coll.h, plus whitebox test updates. No enqueue.cc changes; develop already routes data-movement CE under capture. 1. **Hybrid UC peer writes (ncclPrepUCSync)** - Non-capture: fuse peer WRITE_VALUE + WAIT_VALUE into one batchParams submit (same latency shape as develop). - Capture: issue peer writes as a separate hipStreamBatchMemOp first, then put only waits in batchParams, so the captured stream orders write → wait → reset. Peer values use immediate WRITE_VALUE with GRAPH_SYNC_VALUE while capturing (and the running seq otherwise). We do not use NCCL's D2D memcpy/ceSeqNumDev peer-write path for UC, to avoid a permanent non-graph latency tax. 3. **Even CE intra-batch sync count (ncclCeLaunchBatchOps)** Ported from NCCL: an odd sync count leaves useCompletePtr flipped and desyncs the next ready/complete barrier on replay; append one extra sync when the count would be odd. 4. **Graph-capture window creation** Port NCCL's fix ([22bf876](NVIDIA/nccl@22bf876)) so window creation stays graph-safe during capture. 5. **Tests** Update CeInternalMPITests.cpp for the new ncclPrepUCSync(..., stream) signature (hipStream_t). Non-capture whitebox expectations remain 2*(nRanks-1) ops in the batch. The MC path (`ncclPrepMCSync`) already issues its write separately and is unchanged. ## Issue Tracking JIRA ID: AICOMRCCL-1518 ## Test Plan Validation was performed with and without `NCCL_GRAPH_MIXING_SUPPORT=1`, and with `-G 20`, `-R 2`, and `-x 2`, successfully exercising graph capture and replay for CE collectives. To validate the CE graph path, the current graph guard that prevents RCCL from taking the CE path when graph capture is enabled was temporarily removed. This change was used for validation purposes only and is not included in this PR. The existing graph guard remains in `develop` and is not modified by this PR. CE Graph support will be re-enabled in a follow-up PR once the required HIP stream capture fix #8911 is available in an official ROCm release. At that point, the guard will be converted to ROCm version-based gating, with an environment-variable override if needed. Because the graph guard currently prevents this code path from being exercised in `develop`, no automated **test case is added as part of this PR.** ## Test Result Tested in MI355 with CLR build that include this PR #8911 and in MI450 with latest TheRock nightly build that include this fix. <!-- Pending. --> ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. --------- Co-authored-by: AbandiGa <abandiga@gmail.com>
corey-derochie-amd
pushed a commit
that referenced
this pull request
Oct 15, 2026
## Motivation RCCL's CE data-movement collectives (AllGather, AlltoAll, Scatter, Gather) already run under HIP graph capture on develop, but the UC cross-rank barrier was not safe to replay: peer flag-writes, waits, and capture resets could share one hipStreamBatchMemOp. On HIP that batch does not guarantee wait-before-reset ordering, so a fused reset can clobber a peer flag mid-barrier and hang on graph replay. This PR makes the UC CE barrier graph-replay-safe on HIP, while keeping the fast fused write+wait path for non-capture launches (full NCCL-style always-split peer writes regress small-message latency) ## Technical Details Changes are in ce_coll.cc / ce_coll.h, plus whitebox test updates. No enqueue.cc changes; develop already routes data-movement CE under capture. 1. **Hybrid UC peer writes (ncclPrepUCSync)** - Non-capture: fuse peer WRITE_VALUE + WAIT_VALUE into one batchParams submit (same latency shape as develop). - Capture: issue peer writes as a separate hipStreamBatchMemOp first, then put only waits in batchParams, so the captured stream orders write → wait → reset. Peer values use immediate WRITE_VALUE with GRAPH_SYNC_VALUE while capturing (and the running seq otherwise). We do not use NCCL's D2D memcpy/ceSeqNumDev peer-write path for UC, to avoid a permanent non-graph latency tax. 3. **Even CE intra-batch sync count (ncclCeLaunchBatchOps)** Ported from NCCL: an odd sync count leaves useCompletePtr flipped and desyncs the next ready/complete barrier on replay; append one extra sync when the count would be odd. 4. **Graph-capture window creation** Port NCCL's fix ([22bf876](NVIDIA/nccl@22bf876)) so window creation stays graph-safe during capture. 5. **Tests** Update CeInternalMPITests.cpp for the new ncclPrepUCSync(..., stream) signature (hipStream_t). Non-capture whitebox expectations remain 2*(nRanks-1) ops in the batch. The MC path (`ncclPrepMCSync`) already issues its write separately and is unchanged. ## Issue Tracking JIRA ID: AICOMRCCL-1518 ## Test Plan Validation was performed with and without `NCCL_GRAPH_MIXING_SUPPORT=1`, and with `-G 20`, `-R 2`, and `-x 2`, successfully exercising graph capture and replay for CE collectives. To validate the CE graph path, the current graph guard that prevents RCCL from taking the CE path when graph capture is enabled was temporarily removed. This change was used for validation purposes only and is not included in this PR. The existing graph guard remains in `develop` and is not modified by this PR. CE Graph support will be re-enabled in a follow-up PR once the required HIP stream capture fix #8911 is available in an official ROCm release. At that point, the guard will be converted to ROCm version-based gating, with an environment-variable override if needed. Because the graph guard currently prevents this code path from being exercised in `develop`, no automated **test case is added as part of this PR.** ## Test Result Tested in MI355 with CLR build that include this PR #8911 and in MI450 with latest TheRock nightly build that include this fix. <!-- Pending. --> ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. --------- Co-authored-by: AbandiGa <abandiga@gmail.com>
corey-derochie-amd
pushed a commit
that referenced
this pull request
Oct 15, 2026
## Motivation RCCL's CE data-movement collectives (AllGather, AlltoAll, Scatter, Gather) already run under HIP graph capture on develop, but the UC cross-rank barrier was not safe to replay: peer flag-writes, waits, and capture resets could share one hipStreamBatchMemOp. On HIP that batch does not guarantee wait-before-reset ordering, so a fused reset can clobber a peer flag mid-barrier and hang on graph replay. This PR makes the UC CE barrier graph-replay-safe on HIP, while keeping the fast fused write+wait path for non-capture launches (full NCCL-style always-split peer writes regress small-message latency) ## Technical Details Changes are in ce_coll.cc / ce_coll.h, plus whitebox test updates. No enqueue.cc changes; develop already routes data-movement CE under capture. 1. **Hybrid UC peer writes (ncclPrepUCSync)** - Non-capture: fuse peer WRITE_VALUE + WAIT_VALUE into one batchParams submit (same latency shape as develop). - Capture: issue peer writes as a separate hipStreamBatchMemOp first, then put only waits in batchParams, so the captured stream orders write → wait → reset. Peer values use immediate WRITE_VALUE with GRAPH_SYNC_VALUE while capturing (and the running seq otherwise). We do not use NCCL's D2D memcpy/ceSeqNumDev peer-write path for UC, to avoid a permanent non-graph latency tax. 3. **Even CE intra-batch sync count (ncclCeLaunchBatchOps)** Ported from NCCL: an odd sync count leaves useCompletePtr flipped and desyncs the next ready/complete barrier on replay; append one extra sync when the count would be odd. 4. **Graph-capture window creation** Port NCCL's fix ([22bf876](NVIDIA/nccl@22bf876)) so window creation stays graph-safe during capture. 5. **Tests** Update CeInternalMPITests.cpp for the new ncclPrepUCSync(..., stream) signature (hipStream_t). Non-capture whitebox expectations remain 2*(nRanks-1) ops in the batch. The MC path (`ncclPrepMCSync`) already issues its write separately and is unchanged. ## Issue Tracking JIRA ID: AICOMRCCL-1518 ## Test Plan Validation was performed with and without `NCCL_GRAPH_MIXING_SUPPORT=1`, and with `-G 20`, `-R 2`, and `-x 2`, successfully exercising graph capture and replay for CE collectives. To validate the CE graph path, the current graph guard that prevents RCCL from taking the CE path when graph capture is enabled was temporarily removed. This change was used for validation purposes only and is not included in this PR. The existing graph guard remains in `develop` and is not modified by this PR. CE Graph support will be re-enabled in a follow-up PR once the required HIP stream capture fix #8911 is available in an official ROCm release. At that point, the guard will be converted to ROCm version-based gating, with an environment-variable override if needed. Because the graph guard currently prevents this code path from being exercised in `develop`, no automated **test case is added as part of this PR.** ## Test Result Tested in MI355 with CLR build that include this PR #8911 and in MI450 with latest TheRock nightly build that include this fix. <!-- Pending. --> ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. --------- Co-authored-by: AbandiGa <abandiga@gmail.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.
…tchMemOp
Motivation
Technical Details
Issue Tracking
JIRA ID: AIRUNTIME-2528
Test Plan
Test Result
Submission Checklist