Skip to content

[WIP] Address feedback on ROCM-20519: replace racy pre-sync host read with hipStreamQuery #4256

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Draft
Copilot wants to merge 2 commits into
base: ROCM-20519
Choose a base branch
from
+10 −6
Draft

[WIP] Address feedback on ROCM-20519: replace racy pre-sync host read with hipStreamQuery #4256

Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
ef638f9
Initial plan
Copilot Mar 20, 2026
bc1705d
Replace racy pre-sync host-memory read with hipStreamQuery check
Copilot Mar 20, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
16 changes: 10 additions & 6 deletions projects/hip-tests/catch/unit/device/hipDeviceSynchronize.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -152,12 +152,16 @@ HIP_TEST_CASE(Unit_hipDeviceSynchronize_Functional) {
}


// This check assumes per-stream work is still in flight (typically still in a large H2D
// or the long kernel) when the host reads A[NUM_STREAMS-1][0]. Very fast GPUs or
// HIP_LAUNCH_BLOCKING=true can still fail it; increase copy size via
// HIP_TEST_DEVICE_SYNCHRONIZE_FUNCTIONAL_COPY_MB or drop the assertion if needed.

REQUIRE(NUM_ITERS != A[NUM_STREAMS - 1][0] - 1);
// Verify that work is still in flight on at least one stream without touching shared
// host memory. hipStreamQuery returns hipErrorNotReady when the stream has pending work.
bool anyStreamBusy = false;
for (int i = 0; i < NUM_STREAMS; i++) {
if (hipStreamQuery(stream[i]) == hipErrorNotReady) {
anyStreamBusy = true;
break;
}
}
REQUIRE(anyStreamBusy);
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(NUM_ITERS == A[NUM_STREAMS - 1][0] - 1);
for (int i = 0; i < NUM_STREAMS; i++) {
Expand Down