Repository navigation
fix(hrr): [ROCM-29823] order driver-memcpy D2D against non-blocking readback - #10454
Merged
danielbolsson merged 39 commits intoSep 9, 2026
Merged
Conversation
…eadback
Unit_HRR_DrvMemcpy3D_Direct and Unit_HRR_DrvMemcpy2DUnaligned_Direct read
back a device buffer that nothing guarantees has been written yet, so their
final REQUIRE compares against unwritten memory.
Both workloads issue their last device-to-device copy through the
"synchronous" driver entry point (hipDrvMemcpy3D / hipDrvMemcpy2DUnaligned)
and then read the destination with hipMemcpyAsync on a stream created with
hipStreamNonBlocking. Neither half of that pairing orders the two:
- ihipMemcpyParam3D() deliberately forces the D2D case asynchronous
("Device to Device copies do not need to host side synchronization"),
so the copy is only enqueued on the null stream and the call returns
before it has run.
- A hipStreamNonBlocking stream is excluded from null-stream implicit
synchronization by Device::WaitActiveStreams(), so the readback never
waits for that copy.
The result is a race, not a runtime defect: the runtime behaviour matches
CUDA, which also performs no host-side synchronization for D2D transfers.
Drain the null stream before the readback so the copy is ordered against it.
The hazard was introduced when the 3D workload was rewritten to use a real
copy chain: it previously used hipStreamCreate (a blocking stream, which
does join the null stream) and asserted nothing meaningful. The 2D workload
carries the identical pattern and is fixed here as well.
Measured on gfx950 (MI350X): the readback wins by roughly 12 us of margin,
passing 20/20 below that and failing 20/20 above it, with 15/20 right at the
boundary. That margin is small enough that unrelated timing differences flip
the outcome, which is why the suite is green on gfx950 but fails on gfx1250.
Verified with a standalone reproducer that widens the window with a spin
kernel on the null stream: both workloads read the untouched sentinel before
the change and the copied payload after it.
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. |
andraghetti
approved these changes
Aug 24, 2026
andraghetti
left a comment
Contributor
There was a problem hiding this comment.
Read the diff: the missing hipDeviceSynchronize() before the async D2H readback on a non-blocking stream is the right fix for the gfx1250 race. Happy to see this land before the other HRR PRs.
HIP does not guarantee that hipMalloc returns zeroed memory, so the capture workload can fail before testing replay. Initialize managed memory through uncaptured CPU stores so replay still depends on HRR zero-init while the capture oracle remains deterministic.
Initialize the capture allocation through an unrecorded HSA fill so Linux replay still exercises the ordinary hipMalloc handler while keeping the expected blob deterministic. Native Windows retains the managed-memory fallback because ROCr is unavailable there.
LeoZhao-BJ
reviewed
Aug 27, 2026
The HRR test directory is only added for the AMD platform, and its parent configuration already calls find_package(hip REQUIRED). On non-Windows AMD builds, HIP's package config requires hsa-runtime64 and creates its imported target before this directory is processed. Remove the redundant REQUIRED package lookup here while retaining the direct link needed by Unit_HRR_ZeroInitRead_Direct. Native Windows keeps its existing managed-memory fallback. Co-authored-by: Cursor <cursoragent@cursor.com>
LeoZhao-BJ
approved these changes
Aug 27, 2026
…-memcpy-null-stream-ordering
Unit_HRR_DrvMemcpy3DRoundtrip Unit_HRR_DrvMemcpy2DUnalignedRoundtrip Unit_HRR_ZeroInitRoundtrip
#11157 has landed, so this PR's base auto-retargeted from the migration branch back to develop. Merging develop in replaces the migration content this branch carried from the (now deleted) migration branch with develop's own copy. Two conflicts, both this branch's own additions where develop has the pre-existing form, so both keep this side: * hrr_repro_workload_test.cc — the HSA-vs-Windows split for the never-written source allocation, and the <hsa/hsa_ext_amd.h> include it needs. Develop's ENABLE_ADDRESS_SANITIZER guard around the zeroed-allocation precondition is untouched and still in place. * hrr_workload_test.cc — the two hipDeviceSynchronize() drains after hipDrvMemcpy3D and hipDrvMemcpy2DUnaligned.
Co-authored-by: danielbolsson <113136742+danielbolsson@users.noreply.github.com>
hsa_amd_memory_fill() takes its count in 32-bit words. Passing kN was correct only because the buffer is float[kN]; changing the element type would have silently under- or over-filled it. Use kSZ / sizeof(uint32_t), matching every other call site in the tree.
danielbolsson
deleted the
fix/rocm-29823-drv-memcpy-null-stream-ordering
branch
September 9, 2026 11:59
danielbolsson
added a commit
that referenced
this pull request
Sep 9, 2026
Picks up five commits, including #10454, which landed the driver-memcpy D2D ordering fix in the same test files this stack works on. No conflicts here; the generator reproduces the checked-in artifacts unchanged.
danielbolsson
added a commit
that referenced
this pull request
Sep 9, 2026
… into users/danielbolsson/hrr-regions Carries develop down the stack from the parent branch, picking up #10454. No conflicts; the generator reproduces the checked-in artifacts unchanged.
danielbolsson
added a commit
that referenced
this pull request
Sep 9, 2026
… into users/danielbolsson/hrr-regions-memcpy Carries develop down the stack from the parent branch, picking up #10454. No conflicts; the generator reproduces the checked-in artifacts unchanged.
danielbolsson
added a commit
that referenced
this pull request
Sep 9, 2026
…memcpy' into users/danielbolsson/hrr-test-harness Carries develop down the stack from the parent branch, picking up #10454. No conflicts, despite #10454 having edited hrr_roundtrip_test.cc while this branch restructures the same file. Both landed: #10454's rewritten Unit_HRR_ZeroInitRead_Direct description and its removal of the ASAN guard are present, and the helpers this branch moved into hrr_test_common.hh are still absent from the .cc. Preprocessor nesting is balanced.
danielbolsson
added a commit
that referenced
this pull request
Sep 9, 2026
…ness' into users/danielbolsson/hrr-api-matrix Carries develop down the stack from the parent branch, picking up #10454. No conflicts; the generator reproduces the checked-in artifacts unchanged.
danielbolsson
added a commit
that referenced
this pull request
Sep 9, 2026
…ix' into users/danielbolsson/hrr-api-payloads Carries develop down the stack from the parent branch, picking up #10454. No conflicts. The replay-class overlap guard runs clean and the four classification sets are disjoint, so hipDrvGraphMemcpyNodeSetParams remains in MANUAL_PLAYBACK_APIS only. Artifacts regenerate unchanged.
danielbolsson
added a commit
that referenced
this pull request
Sep 9, 2026
…oads' into users/danielbolsson/hrr-scope-exclusions Carries develop down the stack from the parent branch, picking up #10454. No conflicts; the overlap guard, which this branch extends to cover UNREPLAYABLE_PLAYBACK_APIS, runs clean and the artifacts regenerate unchanged.
danielbolsson
added a commit
that referenced
this pull request
Sep 9, 2026
…clusions' into users/danielbolsson/hrr-kernel-args Carries develop down the stack from the parent branch, picking up #10454. No conflicts; the overlap guard runs clean and the artifacts regenerate unchanged.
Contributor
TheRock Submodule Bump ActivityNewest first
|
danielbolsson
added a commit
that referenced
this pull request
Sep 12, 2026
…10967) Unit_HRR_StreamWriteValue_Direct zeroes its slots with hipMemset and then writes them with hipStreamWriteValue64/32 on a stream created with hipStreamNonBlocking. Nothing orders those two. hipMemset on device memory at offset 0 is asynchronous -- ihipMemset() forces isAsync for it, 'spec says hipMemset will be asynchronous when destination memory is device memory and pointer is non-offseted' (hip_memory.cpp:3444-3454) -- so it is only enqueued on the null stream, and a non-blocking stream never joins the null stream (Device::WaitActiveStreams()). When the zeroing lands after a write the slot reads back 0, which is #10967's '0 == 0xdeadbeeffeedface'. The test already drained after the writes, which cannot help: by then the zeroing may have overwritten them. Drain between the two instead. Same hazard and same fix as the driver-memcpy workloads in #10454; the runtime behaviour is correct and matches CUDA. Add Unit_HRR_NullStreamMemsetOrdering, which pins the contract instead of relying on timing: a spin kernel parked on the null stream ahead of the memset makes an unordered zeroing land after the write on any part. Without the drain it fails 3/3 on gfx950, a part the original test never failed on -- so the failure was never RDNA-specific, only wide enough there to be seen. Co-authored-by: Cursor <cursoragent@cursor.com>
danielbolsson
added a commit
that referenced
this pull request
Sep 14, 2026
…10967) Unit_HRR_StreamWriteValue_Direct zeroes its slots with hipMemset and then writes them with hipStreamWriteValue64/32 on a stream created with hipStreamNonBlocking. Nothing orders those two. hipMemset on device memory at offset 0 is asynchronous -- ihipMemset() forces isAsync for it, 'spec says hipMemset will be asynchronous when destination memory is device memory and pointer is non-offseted' (hip_memory.cpp:3444-3454) -- so it is only enqueued on the null stream, and a non-blocking stream never joins the null stream (Device::WaitActiveStreams()). When the zeroing lands after a write the slot reads back 0, which is #10967's '0 == 0xdeadbeeffeedface'. The test already drained after the writes, which cannot help: by then the zeroing may have overwritten them. Drain between the two instead. Same hazard and same fix as the driver-memcpy workloads in #10454; the runtime behaviour is correct and matches CUDA. Add Unit_HRR_NullStreamMemsetOrdering, which pins the contract instead of relying on timing: a spin kernel parked on the null stream ahead of the memset makes an unordered zeroing land after the write on any part. Without the drain it fails 3/3 on gfx950, a part the original test never failed on -- so the failure was never RDNA-specific, only wide enough there to be seen. Co-authored-by: Cursor <cursoragent@cursor.com>
dlamd1dai
pushed a commit
to dlamd1dai/rocm-systems
that referenced
this pull request
Sep 16, 2026
…eadback (ROCm#10454) ## Motivation `Unit_HRR_DrvMemcpy3D_Direct` and `Unit_HRR_DrvMemcpy2DUnaligned_Direct` read back a device buffer that nothing guarantees has been written yet, so their final `REQUIRE` can compare against unwritten memory. This surfaces as `Unit_HRR_DrvMemcpy3DRoundtrip` failing on gfx1250 with `Capture exit: 42` and a reported mismatch of `0 == 0x5c5c5c5c` — the capture-stage workload failing its own assertion, before playback is ever invoked. The failure is a missing synchronization in the test, not a runtime defect, and it is not gfx1250-specific. ## Technical Details Both workloads issue their last device-to-device copy through the *synchronous* driver entry point (`hipDrvMemcpy3D` / `hipDrvMemcpy2DUnaligned`) and then read the destination with `hipMemcpyAsync` on a stream created with `hipStreamNonBlocking`. Neither half of that pairing orders the two: - `ihipMemcpyParam3D()` deliberately forces the D2D case asynchronous — *"Device to Device copies do not need to host side synchronization"* (`clr/hipamd/src/hip_memory.cpp:2455-2466`). The copy is only enqueued on the null stream and the call returns before it has run. `hipDrvMemcpy2DUnaligned` routes through the same function. - A `hipStreamNonBlocking` stream is excluded from null-stream implicit synchronization by `Device::WaitActiveStreams()` (`clr/hipamd/src/hip_device.cpp:193`), so the readback never waits for that copy. The runtime behaviour is correct and matches CUDA, which likewise performs no host-side synchronization for device-to-device transfers. The fix is on the test side: drain the null stream before the readback. This is the same hazard the HRR playback layer already documents and guards against for its replay zero-init (`clr/hipamd/src/hrr/playback/hip_playback.h:31-44`). **How it was introduced.** The 3D workload previously used `hipStreamCreate` — a *blocking* stream, which does join the null stream — and performed degenerate same-pointer copies whose results were discarded. ROCm#8679 rewrote it to use `hipStreamCreateWithFlags(&s, hipStreamNonBlocking)` with a real A→B→C chain and a live assertion, which is what exposed the ordering gap. The 2D workload carries the identical pattern and is fixed here as well; it is green today only by timing luck. Note that ROCm#9287 is unrelated to this failure despite landing in the same window: `Unit_HRR_DrvMemcpy3DRoundtrip` goes through `hrr_run_exact_roundtrip`, which parses only the pass count and was not touched by that PR, and the failure occurs in `hrr_capture_direct` before playback output is parsed at all. ## Issue Tracking JIRA ID: ROCM-29823/ROCM-30196 ## Test Plan - Full HRR Catch2 suite on a gfx950 (MI350X) host, GPU-free pass and GPU pass. - Standalone reproducer that widens the race window deterministically by parking a spin kernel on the null stream ahead of the driver copy, then reading back on the non-blocking stream. The destination is pre-seeded with an `0xA5A5A5A5` sentinel so an unordered read (sentinel) is distinguishable from data corruption. - Swept the null-stream delay to quantify the margin the unfixed test relies on. - Confirmed both the 3D and 2D paths reproduce the hazard and that the added drain closes both. ## Test Result - **Suite:** GPU-free pass 21/21, GPU pass 71/71, including `Unit_HRR_DrvMemcpy3DRoundtrip` and `Unit_HRR_DrvMemcpy2DUnalignedRoundtrip`. - **Reproducer:** without the fix both `hipDrvMemcpy3D` and `hipDrvMemcpy2DUnaligned` read the untouched sentinel (`h[0]=0xa5a5a5a5`); with the fix both read the copied payload (`h[0]=0x5c5c5c5c`). A blocking stream is ordered either way, matching the analysis. - **Margin:** on gfx950 the readback wins by roughly 12 us — 0/20 raced at 9.1 us of null-stream delay, 15/20 at 13.6 us, 20/20 at 18.2 us and beyond. Partial flakiness at the boundary confirms a race rather than a deterministic behaviour, and a margin that small is easily erased by unrelated timing differences, which is why the suite is green on gfx950 but fails on gfx1250. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. Made with [Cursor](https://cursor.com) --------- Signed-off-by: Daniel Olsson <Daniel.Olsson@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: Mingyang Sun <mingsun@amd.com> Co-authored-by: RahulC <168588218+rahulc-gh@users.noreply.github.com> Co-authored-by: Maneesh Gupta <maneesh.gupta@amd.com> Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com>
anugodavar
pushed a commit
that referenced
this pull request
Sep 17, 2026
…eadback (#10454) ## Motivation `Unit_HRR_DrvMemcpy3D_Direct` and `Unit_HRR_DrvMemcpy2DUnaligned_Direct` read back a device buffer that nothing guarantees has been written yet, so their final `REQUIRE` can compare against unwritten memory. This surfaces as `Unit_HRR_DrvMemcpy3DRoundtrip` failing on gfx1250 with `Capture exit: 42` and a reported mismatch of `0 == 0x5c5c5c5c` — the capture-stage workload failing its own assertion, before playback is ever invoked. The failure is a missing synchronization in the test, not a runtime defect, and it is not gfx1250-specific. ## Technical Details Both workloads issue their last device-to-device copy through the *synchronous* driver entry point (`hipDrvMemcpy3D` / `hipDrvMemcpy2DUnaligned`) and then read the destination with `hipMemcpyAsync` on a stream created with `hipStreamNonBlocking`. Neither half of that pairing orders the two: - `ihipMemcpyParam3D()` deliberately forces the D2D case asynchronous — *"Device to Device copies do not need to host side synchronization"* (`clr/hipamd/src/hip_memory.cpp:2455-2466`). The copy is only enqueued on the null stream and the call returns before it has run. `hipDrvMemcpy2DUnaligned` routes through the same function. - A `hipStreamNonBlocking` stream is excluded from null-stream implicit synchronization by `Device::WaitActiveStreams()` (`clr/hipamd/src/hip_device.cpp:193`), so the readback never waits for that copy. The runtime behaviour is correct and matches CUDA, which likewise performs no host-side synchronization for device-to-device transfers. The fix is on the test side: drain the null stream before the readback. This is the same hazard the HRR playback layer already documents and guards against for its replay zero-init (`clr/hipamd/src/hrr/playback/hip_playback.h:31-44`). **How it was introduced.** The 3D workload previously used `hipStreamCreate` — a *blocking* stream, which does join the null stream — and performed degenerate same-pointer copies whose results were discarded. #8679 rewrote it to use `hipStreamCreateWithFlags(&s, hipStreamNonBlocking)` with a real A→B→C chain and a live assertion, which is what exposed the ordering gap. The 2D workload carries the identical pattern and is fixed here as well; it is green today only by timing luck. Note that #9287 is unrelated to this failure despite landing in the same window: `Unit_HRR_DrvMemcpy3DRoundtrip` goes through `hrr_run_exact_roundtrip`, which parses only the pass count and was not touched by that PR, and the failure occurs in `hrr_capture_direct` before playback output is parsed at all. ## Issue Tracking JIRA ID: ROCM-29823/ROCM-30196 ## Test Plan - Full HRR Catch2 suite on a gfx950 (MI350X) host, GPU-free pass and GPU pass. - Standalone reproducer that widens the race window deterministically by parking a spin kernel on the null stream ahead of the driver copy, then reading back on the non-blocking stream. The destination is pre-seeded with an `0xA5A5A5A5` sentinel so an unordered read (sentinel) is distinguishable from data corruption. - Swept the null-stream delay to quantify the margin the unfixed test relies on. - Confirmed both the 3D and 2D paths reproduce the hazard and that the added drain closes both. ## Test Result - **Suite:** GPU-free pass 21/21, GPU pass 71/71, including `Unit_HRR_DrvMemcpy3DRoundtrip` and `Unit_HRR_DrvMemcpy2DUnalignedRoundtrip`. - **Reproducer:** without the fix both `hipDrvMemcpy3D` and `hipDrvMemcpy2DUnaligned` read the untouched sentinel (`h[0]=0xa5a5a5a5`); with the fix both read the copied payload (`h[0]=0x5c5c5c5c`). A blocking stream is ordered either way, matching the analysis. - **Margin:** on gfx950 the readback wins by roughly 12 us — 0/20 raced at 9.1 us of null-stream delay, 15/20 at 13.6 us, 20/20 at 18.2 us and beyond. Partial flakiness at the boundary confirms a race rather than a deterministic behaviour, and a margin that small is easily erased by unrelated timing differences, which is why the suite is green on gfx950 but fails on gfx1250. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. Made with [Cursor](https://cursor.com) --------- Signed-off-by: Daniel Olsson <Daniel.Olsson@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: Mingyang Sun <mingsun@amd.com> Co-authored-by: RahulC <168588218+rahulc-gh@users.noreply.github.com> Co-authored-by: Maneesh Gupta <maneesh.gupta@amd.com> Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.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.
Motivation
Unit_HRR_DrvMemcpy3D_DirectandUnit_HRR_DrvMemcpy2DUnaligned_Directread back a device buffer that nothing guarantees has been written yet, so their finalREQUIREcan compare against unwritten memory. This surfaces asUnit_HRR_DrvMemcpy3DRoundtripfailing on gfx1250 withCapture exit: 42and a reported mismatch of0 == 0x5c5c5c5c— the capture-stage workload failing its own assertion, before playback is ever invoked.The failure is a missing synchronization in the test, not a runtime defect, and it is not gfx1250-specific.
Technical Details
Both workloads issue their last device-to-device copy through the synchronous driver entry point (
hipDrvMemcpy3D/hipDrvMemcpy2DUnaligned) and then read the destination withhipMemcpyAsyncon a stream created withhipStreamNonBlocking. Neither half of that pairing orders the two:ihipMemcpyParam3D()deliberately forces the D2D case asynchronous — "Device to Device copies do not need to host side synchronization" (clr/hipamd/src/hip_memory.cpp:2455-2466). The copy is only enqueued on the null stream and the call returns before it has run.hipDrvMemcpy2DUnalignedroutes through the same function.hipStreamNonBlockingstream is excluded from null-stream implicit synchronization byDevice::WaitActiveStreams()(clr/hipamd/src/hip_device.cpp:193), so the readback never waits for that copy.The runtime behaviour is correct and matches CUDA, which likewise performs no host-side synchronization for device-to-device transfers. The fix is on the test side: drain the null stream before the readback.
This is the same hazard the HRR playback layer already documents and guards against for its replay zero-init (
clr/hipamd/src/hrr/playback/hip_playback.h:31-44).How it was introduced. The 3D workload previously used
hipStreamCreate— a blocking stream, which does join the null stream — and performed degenerate same-pointer copies whose results were discarded. #8679 rewrote it to usehipStreamCreateWithFlags(&s, hipStreamNonBlocking)with a real A→B→C chain and a live assertion, which is what exposed the ordering gap. The 2D workload carries the identical pattern and is fixed here as well; it is green today only by timing luck.Note that #9287 is unrelated to this failure despite landing in the same window:
Unit_HRR_DrvMemcpy3DRoundtripgoes throughhrr_run_exact_roundtrip, which parses only the pass count and was not touched by that PR, and the failure occurs inhrr_capture_directbefore playback output is parsed at all.Issue Tracking
JIRA ID: ROCM-29823/ROCM-30196
Test Plan
0xA5A5A5A5sentinel so an unordered read (sentinel) is distinguishable from data corruption.Test Result
Unit_HRR_DrvMemcpy3DRoundtripandUnit_HRR_DrvMemcpy2DUnalignedRoundtrip.hipDrvMemcpy3DandhipDrvMemcpy2DUnalignedread the untouched sentinel (h[0]=0xa5a5a5a5); with the fix both read the copied payload (h[0]=0x5c5c5c5c). A blocking stream is ordered either way, matching the analysis.Submission Checklist
Made with Cursor