Repository navigation
fix(clr): wait for a small-batch sync marker after releasing lock - #13532
Merged
Merged
Conversation
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
Contributor
There was a problem hiding this comment.
🔵 Needs a closer look
Signal lifetime and concurrent batch completion need final human review, and the regression test needs consistent device selection.
1 open finding
What changed in this PR
Moves small-batch marker waits outside the queue execution lock so other threads can submit work during stream synchronization.
Changes:
- Separates marker submission from CPU waiting and batch completion.
- Adds a ROCm completion hook while leaving PAL behavior unchanged.
- Adds a watchdog-bounded concurrent stream regression test.
| File | Description |
|---|---|
| projects/hip-tests/catch/unit/stream/hipStreamSyncMultiThread.cc | Adds concurrent synchronization coverage. |
| projects/hip-tests/catch/unit/stream/CMakeLists.txt | Builds the new test. |
| projects/hip-tests/catch/config/configs/unit/stream.yaml | Registers synchronization test tags. |
| projects/clr/rocclr/platform/command.cpp | Calls completion after releasing the lock. |
| projects/clr/rocclr/device/rocm/rocvirtual.hpp | Declares signal-wait and completion interfaces. |
| projects/clr/rocclr/device/rocm/rocvirtual.cpp | Separates fence submission from waiting. |
| projects/clr/rocclr/device/device.hpp | Adds a default no-op completion hook. |
🧠 Review effort: Balanced
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
mangupta
reviewed
Oct 15, 2026
LeoZhao-BJ
force-pushed
the
users/liangzh/ROCM-32517
branch
from
October 16, 2026 02:01
90090bb to
2b18070
Compare
Contributor
|
🤖 AI-assisted pre-review (for @mangupta) Re-reviewed the updated head ( |
saleelk
reviewed
Oct 16, 2026
saleelk
reviewed
Oct 16, 2026
Contributor
|
🤖 AI-assisted pre-review (for @mangupta) Re-reviewed the updated head ( |
Contributor
|
🤖 AI-assisted pre-review (for @mangupta) Re-reviewed the updated head ( |
…e lock HostQueue::finish() sends a Marker with cpu_wait when the pending batch is smaller than DEBUG_CLR_BATCH_CPU_SYNC_SIZE. VirtualGPU::submitMarker() used to flush and busy-wait on the HSA signal inside Command::enqueue(), i.e. while holding the queue's execution() lock. Every other thread that enqueues to the same stream (hipStreamWaitEvent, kernel launches, hipMemcpyAsync, and hipStreamQuery through getLastQueuedCommand) blocked for as long as the GPU took to finish. If the kernel being waited on depends on one of those threads, the process deadlocks, for example torch.distributed gloo block_current_stream. The lock only needs to cover the submission, not the GPU execution, so keep Command::enqueue() and its lock unchanged and move only the wait out of it: - submitMarker() still dispatches the pending fence under the lock, but no longer waits. It attaches the queue's last signal to the marker as its HW event, which retains the signal for the marker's lifetime, so the tracker cannot reuse it in the meantime. - The new VirtualDevice::CompleteCpuWait() is called by Command::enqueue() for such a marker, after it released the lock. The ROCr implementation waits on the signal without the lock, then takes it once to process the signal timing and reset the queue states. The reset now runs only if no work was submitted after the marker, because newer submissions may still use the kernel arguments. The marker's batch is then completed without the lock, since it is private to the marker. The small-batch CPU wait keeps its benefit from #5176 (no extra barrier and no async handler). PAL ignores cpu_wait and keeps its behavior, since the default CompleteCpuWait() is a no-op. A second thread synchronizing the same stream now sees the in-flight marker with a HW event and waits on it outside the lock through the existing IsHwEventReady() path. Add Unit_hipStreamSynchronize_ConcurrentSubmitFromOtherThread. A spin kernel waits on a host flag; while the main thread is in hipStreamSynchronize, other threads poll hipStreamQuery and call hipStreamWaitEvent, and only then release the kernel. A device-side watchdog (wall_clock64 on AMD, clock64 on NVIDIA) turns a deadlock into a test failure. It runs on the null stream and a created stream, with 0 and 3 extra threads synchronizing the same stream. Testing (gfx1201, develop plus rocr-runtime built from the same tree): - the new test fails without the fix (hipStreamQuery and hipStreamWaitEvent blocked until the 10 s watchdog fires) and passes with it - the reproducer from the issue: no async API call blocked for more than 50 ms (167 hipStreamWaitEvent calls before) - several threads synchronizing one stream, with and without concurrent hipStreamWaitEvent/hipStreamQuery: no early return, no deadlock - single-thread finish() latency unchanged - StreamTest and EventTest: no test fails only with the fix Jira: ROCM-32517 GitHub: #11678
- CompleteCpuWait(): check the result of the lock-free wait. If the GPU did not complete the signal, do not wait again under execution() and do not release the queue states, since the work may still be running. The batch is completed as before. - CompleteCpuWait(): document why the batch is completed without the lock. - hipStreamSyncMultiThread: select device 0 before creating resources, and document that the ordering with the main thread is best effort. Jira: ROCM-32517 GitHub: #11678 Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com>
…eCpuWait Address the review comments on CompleteCpuWait(). - Split HwQueueTracker::CpuWaitForSignal() into the wait and the new ProcessSignalTiming(), and drop the wait_only mode. CpuWaitForSignal() is private again and behaves as before for its other callers. - CompleteCpuWait() waits once, without the execution() lock, with WaitForSignal(). It no longer waits a second time under the lock. - WaitForSignal() never returns false. It returns early, with the signal not completed, only if the device is in an error state and HIP_SKIP_ABORT_ON_GPU_ERROR is set. In that case skip the timing processing and the queue state reset, because the GPU may still run, and complete the batch as before, so the threads waiting for its commands don't hang. Only a signal that reached 0 is processed, as IsHwEventReady() does. - Document in CompleteCpuWait() why the batch is completed without the lock. Jira: ROCM-32517 GitHub: #11678 Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com>
litvaOo
force-pushed
the
users/liangzh/ROCM-32517
branch
from
October 19, 2026 14:49
abcdc14 to
611bfff
Compare
Contributor
|
🤖 AI-assisted pre-review (for @mangupta) Re-reviewed the updated head ( |
LeoZhao-BJ
force-pushed
the
users/liangzh/ROCM-32517
branch
from
October 21, 2026 02:15
9f6875a to
611bfff
Compare
Contributor
|
🤖 AI-assisted pre-review (for @mangupta) Re-reviewed the updated head ( |
gandryey
approved these changes
Oct 22, 2026
Contributor
TheRock Submodule Bump ActivityNewest first
|
saleelk
added a commit
that referenced
this pull request
Oct 27, 2026
…(part 2) Part 2 of a series moving CLR completion signals to the ROCr batch API and device memory. This part covers the queue completion signals. DEBUG_CLR_DEV_SIGNALS (default off) places each queue's completion signals in device memory when the device has a large BAR and no profiler tool is attached, the same condition as the graph signals, in the same pool (uncached on gfx94x/gfx95x, coarse-grained on gfx125x). Every host access to those signals then crosses PCIe, so VirtualGPU uses a new RetiringSignalTracker, which proves that a signal's last use finished from host-side state instead of reading the signal, and re-arms signals ahead of use in batches. Otherwise VirtualGPU keeps LegacySignalTracker. Device-memory signals, graph and queue, are HIP-only: OpenCL keeps host signals and LegacySignalTracker. RetiringSignalTracker: - Retirement. Each use records its AQL slot, whether the packet had the barrier bit, and the HW queue generation. Barrier-bit packets are recorded as fences (NoteFence): once the queue read index passes a fence, every earlier packet has completed. A completed fence use proves everything up to and including it, and reading the newest fence in a batch (ProbeFence) can retire the whole batch with one load. A HW queue change (NoteQueueChange) starts a new generation. - Arming. Reclaim walks the ring from the oldest use, re-arms retired slots with direct stores and one HDP flush per 32 slots, and stops at the first busy one. A slot whose signal is still shared (an event, a GPU wait or a pending handler) gets an armed spare instead; the old signal is parked and reused once the tracker holds the last reference. A ring about to reuse a busy slot grows, up to 16 x ROC_SIGNAL_POOL_SIZE, before waiting on the oldest signal. - GPU waits. Signals returned by WaitingSignal for a GPU waiter are pinned with a reference until the consuming packets have passed and every earlier use has retired, so re-arming cannot hide the old use from the waiter. A drained queue (NoteQueueDrained) releases them. - HDP flushes. When the device's AQL rings are in device memory, the arming stores are not flushed: the packets that use the signals reach the ring through the same BAR after the stores, in order, and a waiter on another device loads the value before it waits. A signal handed to a ROCr copy is flushed when its use begins, since ROCr's SDMA and blit rings are in system memory. - The async handler holds a reference to the signal until it has processed it, and drops it before the VirtualGPU reference and inside the async-handler scope, so neither the signal nor the HW queue is destroyed on the handler thread outside the deferral path. Shared changes: - ActiveSignal takes the AQL slot and barrier bit of the packet the signal completes; dispatchBarrierValuePacket reserves its slot before taking the signal. - HwQueueTracker::Create builds the completion-signal ring (64) and the interrupt pool (5) as two batches for both trackers, in device memory with DEBUG_CLR_DEV_SIGNALS. - The kernarg and staging ManagedBuffer chunk signals are a batch too, in device memory with DEBUG_CLR_DEV_SIGNALS. A chunk signal there is re-armed with a direct store, flushed only when the AQL ring is in system memory, since a store through ROCr reads it back across PCIe. - Timing of a use known to be complete (after a CPU wait, or proven retired) skips the signal value check in CacheTimingData. - A ROCr submission that fails after issuing some of the sub-operations sharing one signal (DMA copyBufferRect line by line, SVM prefetch batch) waits for the issued ones (WaitForIssuedOps) before resetting the slot, with either tracker: their late decrements would otherwise complete the slot's next use early. - WaitForSignal no longer loads the signal before waiting, since the wait loads it first, and a blocking wait returns as soon as its 100 us active phase sees completion instead of also running the blocked wait. - The small-batch sync path (#13532) reads the marker's signal once: CompleteCpuWait checks the GPU-error state directly instead of reloading the signal, and HostQueue::finish skips the HW-event wait for a marker that enqueue already completed. Test (MI300X): - HIP stream/graph test with an LD_PRELOAD batch counter passes with DEBUG_CLR_DEV_SIGNALS=0 and 1; graph_bench --verify passes. - With DEBUG_CLR_DEV_SIGNALS=1, median hipEventSynchronize after an empty kernel drops from 7.2 to 4.7 us (10.9 to 4.7 us for blocking-sync events), and hipStreamSynchronize on completed work from 7.5 to 6.1 us; host-memory signals are unchanged. - Unit_hipStreamSynchronize_ConcurrentSubmitFromOtherThread (#13532) passes with the flag off and on, with DEBUG_CLR_BATCH_CPU_SYNC_SIZE at its default and at 0. - hip-tests StreamTest, GraphsTest1, GraphsTest2: failure sets identical to part 1 with the flag off and on (all pre-existing). - rocvirtual.cpp, rocdevice.cpp and rocblit.cpp compile with ROCR_DYN_DLL defined.
saleelk
added a commit
that referenced
this pull request
Oct 27, 2026
…(part 2) Part 2 of a series moving CLR completion signals to the ROCr batch API and device memory. This part covers the queue completion signals. DEBUG_CLR_DEV_SIGNALS (default off) places each queue's completion signals in device memory when the device has a large BAR and no profiler tool is attached, the same condition as the graph signals, in the same pool (uncached on gfx94x/gfx95x, coarse-grained on gfx125x). Every host access to those signals then crosses PCIe, so VirtualGPU uses a new RetiringSignalTracker, which proves that a signal's last use finished from host-side state instead of reading the signal, and re-arms signals ahead of use in batches. Otherwise VirtualGPU keeps LegacySignalTracker. Device-memory signals, graph and queue, are HIP-only: OpenCL keeps host signals and LegacySignalTracker. RetiringSignalTracker: - Retirement. Each use records its AQL slot, whether the packet had the barrier bit, and the HW queue generation. Barrier-bit packets are recorded as fences (NoteFence): once the queue read index passes a fence, every earlier packet has completed. A completed fence use proves everything up to and including it, and reading the newest fence in a batch (ProbeFence) can retire the whole batch with one load. A HW queue change (NoteQueueChange) starts a new generation. - Arming. Reclaim walks the ring from the oldest use, re-arms retired slots with direct stores and one HDP flush per 32 slots, and stops at the first busy one. A slot whose signal is still shared (an event, a GPU wait or a pending handler) gets an armed spare instead; the old signal is parked and reused once the tracker holds the last reference. A ring about to reuse a busy slot grows, up to 16 x ROC_SIGNAL_POOL_SIZE, before waiting on the oldest signal. - Timing. A use's timing is collected before its signal is re-armed, so no Timestamp lists an armed signal. A command whose signals wrap the ring, such as a graph launch under the built-in profiler, would otherwise have its async handler wait on a re-armed signal while holding the Timestamp the next use needs, and read the next use's timestamps. - GPU waits. Signals returned by WaitingSignal for a GPU waiter are pinned with a reference until the consuming packets have passed and every earlier use has retired, so re-arming cannot hide the old use from the waiter. A drained queue (NoteQueueDrained) releases them. - HDP flushes. When the device's AQL rings are in device memory, the arming stores are not flushed: the packets that use the signals reach the ring through the same BAR after the stores, in order, and a waiter on another device loads the value before it waits. A signal handed to a ROCr copy is flushed when its use begins, since ROCr's SDMA and blit rings are in system memory. - The async handler holds a reference to the signal until it has processed it, and drops it before the VirtualGPU reference and inside the async-handler scope, so neither the signal nor the HW queue is destroyed on the handler thread outside the deferral path. Shared changes: - ActiveSignal takes the AQL slot and barrier bit of the packet the signal completes; dispatchBarrierValuePacket reserves its slot before taking the signal. - HwQueueTracker::Create builds the completion-signal ring (64) and the interrupt pool (5) as two batches for both trackers, in device memory with DEBUG_CLR_DEV_SIGNALS. - The kernarg and staging ManagedBuffer chunk signals are a batch too, in device memory with DEBUG_CLR_DEV_SIGNALS. A chunk signal there is re-armed with a direct store, flushed only when the AQL ring is in system memory, since a store through ROCr reads it back across PCIe. - Timing of a use known to be complete (after a CPU wait, or proven retired) skips the signal value check in CacheTimingData. - A ROCr submission that fails after issuing some of the sub-operations sharing one signal (DMA copyBufferRect line by line, SVM prefetch batch) waits for the issued ones (WaitForIssuedOps) before resetting the slot, with either tracker: their late decrements would otherwise complete the slot's next use early. - WaitForSignal no longer loads the signal before waiting, since the wait loads it first, and a blocking wait returns as soon as its 100 us active phase sees completion instead of also running the blocked wait. - The small-batch sync path (#13532) reads the marker's signal once: CompleteCpuWait checks the GPU-error state directly instead of reloading the signal, and HostQueue::finish skips the HW-event wait for a marker that enqueue already completed. Test (MI300X): - HIP stream/graph test with an LD_PRELOAD batch counter passes with DEBUG_CLR_DEV_SIGNALS=0 and 1; graph_bench --verify passes. - With DEBUG_CLR_DEV_SIGNALS=1, median hipEventSynchronize after an empty kernel drops from 7.2 to 4.7 us (10.9 to 4.7 us for blocking-sync events), and hipStreamSynchronize on completed work from 7.5 to 6.1 us; host-memory signals are unchanged. - Unit_hipStreamSynchronize_ConcurrentSubmitFromOtherThread (#13532) passes with the flag off and on, with DEBUG_CLR_BATCH_CPU_SYNC_SIZE at its default and at 0. - hip_launch_latency -k -g, plain and with GPU_CLR_PROFILE_OUTPUT (a signal per dispatch), completes with --verify counts correct, with the flag off and on. - hip-tests StreamTest, GraphsTest1, GraphsTest2: failure sets identical to part 1 with the flag off and on (all pre-existing). - rocvirtual.cpp, rocdevice.cpp and rocblit.cpp compile with ROCR_DYN_DLL defined.
saleelk
added a commit
that referenced
this pull request
Oct 27, 2026
…(part 2) Part 2 of a series moving CLR completion signals to the ROCr batch API and device memory. This part covers the queue completion signals. DEBUG_CLR_DEV_SIGNALS (default off) places each queue's completion signals in device memory when the device has a large BAR and no profiler tool is attached, the same condition as the graph signals, in the same pool (uncached on gfx94x/gfx95x, coarse-grained on gfx125x). Every host access to those signals then crosses PCIe, so VirtualGPU uses a new RetiringSignalTracker, which proves that a signal's last use finished from host-side state instead of reading the signal, and re-arms signals ahead of use in batches. Otherwise VirtualGPU keeps LegacySignalTracker. Device-memory signals, graph and queue, are HIP-only: OpenCL keeps host signals and LegacySignalTracker. RetiringSignalTracker: - Retirement. Each use records its AQL slot, whether the packet had the barrier bit, and the HW queue generation. Barrier-bit packets are recorded as fences (NoteFence): once the queue read index passes a fence, every earlier packet has completed. A completed fence use proves everything up to and including it, and reading the newest fence in a batch (ProbeFence) can retire the whole batch with one load. A HW queue change (NoteQueueChange) starts a new generation. - Arming. Reclaim walks the ring from the oldest use, re-arms retired slots with direct stores and one HDP flush per 32 slots, and stops at the first busy one. A slot whose signal is still shared (an event, a GPU wait or a pending handler) gets an armed spare instead; the old signal is parked and reused once the tracker holds the last reference. A ring about to reuse a busy slot grows, up to 16 x ROC_SIGNAL_POOL_SIZE, before waiting on the oldest signal. - Timing. A use's timing is collected before its signal is re-armed, so no Timestamp lists an armed signal. A command whose signals wrap the ring, such as a graph launch under the built-in profiler, would otherwise have its async handler wait on a re-armed signal while holding the Timestamp the next use needs, and read the next use's timestamps. - Timing-only uses. A dispatch that takes a signal only because the built-in profiler times it gets an armed host-memory signal from a separate pool (PlaceForUse), so the handler reads its value and timestamps without crossing PCIe. A use that a GPU waiter or a sync may read swaps a device-memory signal back into the slot. - GPU waits. Signals returned by WaitingSignal for a GPU waiter are pinned with a reference until the consuming packets have passed and every earlier use has retired, so re-arming cannot hide the old use from the waiter. A drained queue (NoteQueueDrained) releases them. - HDP flushes. When the device's AQL rings are in device memory, the arming stores are not flushed: the packets that use the signals reach the ring through the same BAR after the stores, in order, and a waiter on another device loads the value before it waits. A signal handed to a ROCr copy is flushed when its use begins, since ROCr's SDMA and blit rings are in system memory. - The async handler holds a reference to the signal until it has processed it, and drops it before the VirtualGPU reference and inside the async-handler scope, so neither the signal nor the HW queue is destroyed on the handler thread outside the deferral path. Shared changes: - ActiveSignal takes the AQL slot and barrier bit of the packet the signal completes; dispatchBarrierValuePacket reserves its slot before taking the signal. - HwQueueTracker::Create builds the completion-signal ring (64) and the interrupt pool (5) as two batches for both trackers, in device memory with DEBUG_CLR_DEV_SIGNALS. - The kernarg and staging ManagedBuffer chunk signals are a batch too, in device memory with DEBUG_CLR_DEV_SIGNALS. A chunk signal there is re-armed with a direct store, flushed only when the AQL ring is in system memory, since a store through ROCr reads it back across PCIe. - Timing of a use known to be complete (after a CPU wait, or proven retired) skips the signal value check in CacheTimingData. - A ROCr submission that fails after issuing some of the sub-operations sharing one signal (DMA copyBufferRect line by line, SVM prefetch batch) waits for the issued ones (WaitForIssuedOps) before resetting the slot, with either tracker: their late decrements would otherwise complete the slot's next use early. - WaitForSignal no longer loads the signal before waiting, since the wait loads it first, and a blocking wait returns as soon as its 100 us active phase sees completion instead of also running the blocked wait. - hipEventSynchronize completes the queue's submission batch through the event's marker once the marker is done (HostQueue::CompleteSubmissionBatch). A record marker doesn't own a batch, so in a record/sync loop every command, and its signal, stayed referenced until the batch reached DEBUG_CLR_MAX_BATCH_SIZE. With device signals the tracker parked those signals and refilled its pools with 64-signal allocations inside the timed intervals. - The small-batch sync path (#13532) reads the marker's signal once: CompleteCpuWait checks the GPU-error state directly instead of reloading the signal, and HostQueue::finish skips the HW-event wait for a marker that enqueue already completed. Test (MI300X): - HIP stream/graph test with an LD_PRELOAD batch counter passes with DEBUG_CLR_DEV_SIGNALS=0 and 1; graph_bench --verify passes. - With DEBUG_CLR_DEV_SIGNALS=1, median hipEventSynchronize after an empty kernel drops from 7.2 to 4.7 us (10.9 to 4.7 us for blocking-sync events), and hipStreamSynchronize on completed work from 7.5 to 6.1 us; host-memory signals are unchanged. - Unit_hipStreamSynchronize_ConcurrentSubmitFromOtherThread (#13532) passes with the flag off and on, with DEBUG_CLR_BATCH_CPU_SYNC_SIZE at its default and at 0. - hip_launch_latency -k -g, plain and with GPU_CLR_PROFILE_OUTPUT (a signal per dispatch), completes with --verify counts correct, with the flag off and on. With GPU_CLR_PROFILE_OUTPUT and the flag on, graph enqueue takes 32-42 us (44-47 us with host signals) and a single launch measured with events 5.3-5.5 us (6.9 us). - A hipEventRecord/kernel/hipEventRecord/hipEventSynchronize loop no longer parks signals; its median event time with the flag on is 3.5-3.7 us (5.4 us with host signals), without the earlier 75-575 us outliers. - hip-tests EventTest, StreamTest, GraphsTest1, GraphsTest2: failure sets identical with the flag off and on, and to a build without the timing-only signals and the hipEventSynchronize batch completion (all pre-existing: order-dependent sticky errors, and a stream capture a graph test leaves open). - rocvirtual.cpp, rocdevice.cpp and rocblit.cpp compile with ROCR_DYN_DLL defined.
saleelk
added a commit
that referenced
this pull request
Oct 27, 2026
…(part 2) Part 2 of a series moving CLR completion signals to the ROCr batch API and device memory. This part covers the queue completion signals. DEBUG_CLR_DEV_SIGNALS (default off) places each queue's completion signals in device memory when the device has a large BAR and no profiler tool is attached, the same condition as the graph signals, in the same pool (uncached on gfx94x/gfx95x, coarse-grained on gfx125x). Every host access to those signals then crosses PCIe, so VirtualGPU uses a new RetiringSignalTracker, which proves that a signal's last use finished from host-side state instead of reading the signal, and re-arms signals ahead of use in batches. Otherwise VirtualGPU keeps LegacySignalTracker. Device-memory signals, graph and queue, are HIP-only: OpenCL keeps host signals and LegacySignalTracker. RetiringSignalTracker: - Retirement. Each use records its AQL slot, whether the packet had the barrier bit, and the HW queue generation. Barrier-bit packets are recorded as fences (NoteFence): once the queue read index passes a fence, every earlier packet has completed. A completed fence use proves everything up to and including it, and reading the newest fence in a batch (ProbeFence) can retire the whole batch with one load. A HW queue change (NoteQueueChange) starts a new generation. - Arming. Reclaim walks the ring from the oldest use, re-arms retired slots with direct stores and one HDP flush per 32 slots, and stops at the first busy one. A slot whose signal is still shared (an event, a GPU wait or a pending handler) gets an armed spare instead; the old signal is parked and reused once the tracker holds the last reference. A ring about to reuse a busy slot grows, up to 16 x ROC_SIGNAL_POOL_SIZE, before waiting on the oldest signal. - Timing. A use's timing is collected before its signal is re-armed, so no Timestamp lists an armed signal. A command whose signals wrap the ring, such as a graph launch under the built-in profiler, would otherwise have its async handler wait on a re-armed signal while holding the Timestamp the next use needs, and read the next use's timestamps. - Timing-only uses. A dispatch that takes a signal only because the built-in profiler times it gets an armed host-memory signal from a separate pool (PlaceForUse), so the handler reads its value and timestamps without crossing PCIe. A use that a GPU waiter or a sync may read swaps a device-memory signal back into the slot. A slot keeps its signal when the replacement's allocation fell back to host memory. - GPU waits. Signals returned by WaitingSignal for a GPU waiter are pinned with a reference until the consuming packets have passed and every earlier use has retired, so re-arming cannot hide the old use from the waiter. A drained queue (NoteQueueDrained) releases them. - HDP flushes. When the device's AQL rings are in device memory, the arming stores are not flushed: the packets that use the signals reach the ring through the same BAR after the stores, in order, and a waiter on another device loads the value before it waits. A signal handed to a ROCr copy is flushed when its use begins, since ROCr's SDMA and blit rings are in system memory. - The async handler holds a reference to the signal until it has processed it, and drops it before the VirtualGPU reference and inside the async-handler scope, so neither the signal nor the HW queue is destroyed on the handler thread outside the deferral path. Signal allocation, with DEBUG_CLR_DEV_SIGNALS only. With the flag off, every queue and ManagedBuffer signal is created with hsa_signal_create as before: - HwQueueTracker::Create builds the completion-signal ring (64) and the interrupt pool (5) as two device-memory batches. - The kernarg and staging ManagedBuffer chunk signals are a device-memory batch too. A chunk signal there is re-armed with a direct store, flushed only when the AQL ring is in system memory, since a store through ROCr reads it back across PCIe. Shared changes. These apply with DEBUG_CLR_DEV_SIGNALS off too, so turning the flag off does not restore the previous behavior in these paths: - ActiveSignal takes the AQL slot and barrier bit of the packet the signal completes; dispatchBarrierValuePacket reserves its slot before taking the signal. - Timing of a use known to be complete (after a CPU wait, or proven retired) skips the signal value check in CacheTimingData. - A ROCr submission that fails after issuing some of the sub-operations sharing one signal (DMA copyBufferRect line by line, SVM prefetch batch) waits for the issued ones (WaitForIssuedOps) before resetting the slot, with either tracker: their late decrements would otherwise complete the slot's next use early. - WaitForSignal no longer loads the signal before waiting, since the wait loads it first, and a blocking wait returns as soon as its 100 us active phase sees completion instead of also running the blocked wait. - hipEventSynchronize completes the queue's submission batch through the event's marker once the marker is done (HostQueue::CompleteSubmissionBatch). A record marker doesn't own a batch, so in a record/sync loop every command, and its signal, stayed referenced until the batch reached DEBUG_CLR_MAX_BATCH_SIZE. With device signals the tracker parked those signals and refilled its pools with 64-signal allocations inside the timed intervals. The call returns without the queue lock when the marker already completed. This is the same batch completion hipStreamSynchronize runs on the calling thread. - The small-batch sync path (#13532) reads the marker's signal once: CompleteCpuWait checks the GPU-error state directly instead of reloading the signal, and HostQueue::finish skips the HW-event wait for a marker that enqueue already completed. Test (MI300X): - HIP stream/graph test with an LD_PRELOAD batch counter passes with DEBUG_CLR_DEV_SIGNALS=0 and 1; graph_bench --verify passes. - With DEBUG_CLR_DEV_SIGNALS=1, median hipEventSynchronize after an empty kernel drops from 7.2 to 4.7 us (10.9 to 4.7 us for blocking-sync events), and hipStreamSynchronize on completed work from 7.5 to 6.1 us; host-memory signals are unchanged. - Unit_hipStreamSynchronize_ConcurrentSubmitFromOtherThread (#13532) passes with the flag off and on, with DEBUG_CLR_BATCH_CPU_SYNC_SIZE at its default and at 0. - hip_launch_latency -k -g, plain and with GPU_CLR_PROFILE_OUTPUT (a signal per dispatch), completes with --verify counts correct, with the flag off and on. With GPU_CLR_PROFILE_OUTPUT and the flag on, graph enqueue takes 32-42 us (44-47 us with host signals) and a single launch measured with events 5.3-5.5 us (6.9 us). - A hipEventRecord/kernel/hipEventRecord/hipEventSynchronize loop no longer parks signals; its median event time with the flag on is 3.5-3.7 us (5.4 us with host signals), without the earlier 75-575 us outliers. - hip-tests EventTest, StreamTest, GraphsTest1, GraphsTest2: failure sets identical with the flag off and on, and to a build without the timing-only signals and the hipEventSynchronize batch completion (all pre-existing: order-dependent sticky errors, and a stream capture a graph test leaves open). - With DEBUG_CLR_DEV_SIGNALS=0 no signal goes through hsa_amd_signal_batch_create (checked with an LD_PRELOAD counter). - With every device-memory batch failing (an LD_PRELOAD shim), a 3000-iteration event loop allocates one spare batch and keeps it; before, it allocated 98 batches and grew the host pool to about 6000 signals. - rocvirtual.cpp, rocdevice.cpp and rocblit.cpp compile with ROCR_DYN_DLL defined. JIRA ID : AIRUNTIME-2943
saleelk
added a commit
that referenced
this pull request
Oct 28, 2026
…(part 2) Part 2 of a series moving CLR completion signals to the ROCr batch API and device memory. This part covers the queue completion signals. DEBUG_CLR_DEV_SIGNALS (default off) places each queue's completion signals in device memory when the device has a large BAR and no profiler tool is attached, the same condition as the graph signals, in the same pool (uncached on gfx94x/gfx95x, coarse-grained on gfx125x). Every host access to those signals then crosses PCIe, so VirtualGPU uses a new RetiringSignalTracker, which proves that a signal's last use finished from host-side state instead of reading the signal, and re-arms signals ahead of use in batches. Otherwise VirtualGPU keeps LegacySignalTracker. Device-memory signals, graph and queue, are HIP-only: OpenCL keeps host signals and LegacySignalTracker. RetiringSignalTracker: - Retirement. Each use records its AQL slot, whether the packet had the barrier bit, and the HW queue generation. A use published with the barrier bit is a fence: its completion proves that every earlier packet on the HW queue completed. The packet processor may advance the read index before a packet completes (HSA System Architecture 2.8.3), so the read index only nominates the newest fence it passed, and one load of that fence's signal confirms it. Reading the newest fence in a batch (ProbeFence) likewise retires the whole batch with one load. A HW queue change (NoteQueueChange) starts a new generation. - Arming. Reclaim walks the ring from the oldest use, re-arms retired slots with direct stores and one HDP flush per 32 slots, and stops at the first busy one. A slot whose signal is still shared (an event, a GPU wait or a pending handler) gets an armed spare instead; the old signal is parked and reused once the tracker holds the last reference. A ring about to reuse a busy slot grows, up to 16 x ROC_SIGNAL_POOL_SIZE, before waiting on the oldest signal. - Timing. A use's timing is collected before its signal is re-armed, so no Timestamp lists an armed signal. A command whose signals wrap the ring, such as a graph launch under the built-in profiler, would otherwise have its async handler wait on a re-armed signal while holding the Timestamp the next use needs, and read the next use's timestamps. - Timing-only uses. A dispatch that takes a signal only because the built-in profiler times it gets an armed host-memory signal from a separate pool (PlaceForUse), so the handler reads its value and timestamps without crossing PCIe. A use that a GPU waiter or a sync may read swaps a device-memory signal back into the slot. A slot keeps its signal when the replacement's allocation fell back to host memory. - GPU waits. Signals returned by WaitingSignal for a GPU waiter are pinned with a reference until the consuming packets have passed and every earlier use has retired, so re-arming cannot hide the old use from the waiter. A drained queue (NoteQueueDrained) releases them. - HDP flushes. When the device's AQL rings are in device memory, the arming stores are not flushed: the packets that use the signals reach the ring through the same BAR after the stores, in order, and a waiter on another device loads the value before it waits. A signal handed to a ROCr copy is flushed when its use begins, since ROCr's SDMA and blit rings are in system memory. - The async handler holds a reference to the signal until it has processed it, and drops it before the VirtualGPU reference and inside the async-handler scope, so neither the signal nor the HW queue is destroyed on the handler thread outside the deferral path. Signal allocation, with DEBUG_CLR_DEV_SIGNALS only. With the flag off, every queue and ManagedBuffer signal is created with hsa_signal_create as before: - HwQueueTracker::Create builds the completion-signal ring (64) and the interrupt pool (5) as two device-memory batches. - The kernarg and staging ManagedBuffer chunk signals are a device-memory batch too. A chunk signal there is re-armed with a direct store, flushed only when the AQL ring is in system memory, since a store through ROCr reads it back across PCIe. Shared changes. These apply with DEBUG_CLR_DEV_SIGNALS off too, so turning the flag off does not restore the previous behavior in these paths: - ActiveSignal takes the AQL slot and barrier bit of the packet the signal completes; dispatchBarrierValuePacket reserves its slot before taking the signal. - Timing of a use known to be complete (after a CPU wait, or proven retired) skips the signal value check in CacheTimingData. - A ROCr submission that fails after issuing some of the sub-operations sharing one signal (DMA copyBufferRect line by line, SVM prefetch batch) waits for the issued ones (WaitForIssuedOps) before resetting the slot, with either tracker: their late decrements would otherwise complete the slot's next use early. - WaitForSignal no longer loads the signal before waiting, since the wait loads it first, and a blocking wait returns as soon as its 100 us active phase sees completion instead of also running the blocked wait. - hipEventSynchronize completes the queue's submission batch through the event's marker once the marker is done (HostQueue::CompleteSubmissionBatch). A record marker doesn't own a batch, so in a record/sync loop every command, and its signal, stayed referenced until the batch reached DEBUG_CLR_MAX_BATCH_SIZE. With device signals the tracker parked those signals and refilled its pools with 64-signal allocations inside the timed intervals. The call returns without the queue lock when the marker already completed. This is the same batch completion hipStreamSynchronize runs on the calling thread. - The small-batch sync path (#13532) reads the marker's signal once: CompleteCpuWait checks the GPU-error state directly instead of reloading the signal, and HostQueue::finish skips the HW-event wait for a marker that enqueue already completed. Test (MI300X): - HIP stream/graph test with an LD_PRELOAD batch counter passes with DEBUG_CLR_DEV_SIGNALS=0 and 1; graph_bench --verify passes. - With DEBUG_CLR_DEV_SIGNALS=1, median hipEventSynchronize after an empty kernel drops from 7.2 to 4.7 us (10.9 to 4.7 us for blocking-sync events), and hipStreamSynchronize on completed work from 7.5 to 6.1 us; host-memory signals are unchanged. - Unit_hipStreamSynchronize_ConcurrentSubmitFromOtherThread (#13532) passes with the flag off and on, with DEBUG_CLR_BATCH_CPU_SYNC_SIZE at its default and at 0. - hip_launch_latency -k -g, plain and with GPU_CLR_PROFILE_OUTPUT (a signal per dispatch), completes with --verify counts correct, with the flag off and on. With GPU_CLR_PROFILE_OUTPUT and the flag on, graph enqueue takes 32-42 us (44-47 us with host signals) and a single launch measured with events 5.3-5.5 us (6.9 us). - A hipEventRecord/kernel/hipEventRecord/hipEventSynchronize loop no longer parks signals; its median event time with the flag on is 3.5-3.7 us (5.4 us with host signals), without the earlier 75-575 us outliers. - Confirming the fence costs nothing measurable: hip_launch_latency -k -g and a 12-stream test sharing 4 HW queues (6.3-6.4 us per launch) match the read-index-only version. On MI300X the read index never passed an incomplete packet with a spinning kernel ahead of a barrier-AND, barrier-value or kernel fence. - hip-tests EventTest, StreamTest, GraphsTest1, GraphsTest2: failure sets identical with the flag off and on, and to a build without the timing-only signals and the hipEventSynchronize batch completion (all pre-existing: order-dependent sticky errors, and a stream capture a graph test leaves open). - With DEBUG_CLR_DEV_SIGNALS=0 no signal goes through hsa_amd_signal_batch_create (checked with an LD_PRELOAD counter). - With every device-memory batch failing (an LD_PRELOAD shim), a 3000-iteration event loop allocates one spare batch and keeps it; before, it allocated 98 batches and grew the host pool to about 6000 signals. - rocvirtual.cpp, rocdevice.cpp and rocblit.cpp compile with ROCR_DYN_DLL defined. JIRA ID : AIRUNTIME-2943
saleelk
added a commit
that referenced
this pull request
Oct 28, 2026
…(part 2) Part 2 of a series moving CLR completion signals to the ROCr batch API and device memory. This part covers the queue completion signals. DEBUG_CLR_DEV_SIGNALS (default off) places each queue's completion signals in device memory when the device has a large BAR and no profiler tool is attached, the same condition as the graph signals, in the same pool (uncached on gfx94x/gfx95x, coarse-grained on gfx125x). Every host access to those signals then crosses PCIe, so VirtualGPU uses a new RetiringSignalTracker, which proves that a signal's last use finished from host-side state instead of reading the signal, and re-arms signals ahead of use in batches. Otherwise VirtualGPU keeps LegacySignalTracker. Device-memory signals, graph and queue, are HIP-only: OpenCL keeps host signals and LegacySignalTracker. RetiringSignalTracker: - Retirement. Each use records its AQL slot, whether the packet had the barrier bit, and the HW queue generation. A use published with the barrier bit is a fence: its completion proves that every earlier packet on the HW queue completed. The packet processor may advance the read index before a packet completes (HSA System Architecture 2.8.3), so the read index only nominates the newest fence it passed, and one load of that fence's signal confirms it. Reading the newest fence in a batch (ProbeFence) likewise retires the whole batch with one load. A HW queue change (NoteQueueChange) starts a new generation. - Arming. Reclaim walks the ring from the oldest use, re-arms retired slots with direct stores and one HDP flush per 32 slots, and stops at the first busy one. A slot whose signal is still shared (an event, a GPU wait or a pending handler) gets an armed spare instead; the old signal is parked and reused once the tracker holds the last reference. A ring about to reuse a busy slot grows, up to 16 x ROC_SIGNAL_POOL_SIZE, before waiting on the oldest signal. - Timing. A use's timing is collected before its signal is re-armed, so no Timestamp lists an armed signal. A command whose signals wrap the ring, such as a graph launch under the built-in profiler, would otherwise have its async handler wait on a re-armed signal while holding the Timestamp the next use needs, and read the next use's timestamps. - Timing-only uses. A dispatch that takes a signal only because the built-in profiler times it gets an armed host-memory signal from a separate pool (PlaceForUse), so the handler reads its value and timestamps without crossing PCIe. A use that a GPU waiter or a sync may read swaps a device-memory signal back into the slot. A slot keeps its signal when the replacement's allocation fell back to host memory. - GPU waits. Signals returned by WaitingSignal for a GPU waiter are pinned with a reference until the consuming packets have passed and every earlier use has retired, so re-arming cannot hide the old use from the waiter. A drained queue (NoteQueueDrained) releases them. - HDP flushes. When the device's AQL rings are in device memory, the arming stores are not flushed: the packets that use the signals reach the ring through the same BAR after the stores, in order, and a waiter on another device loads the value before it waits. A signal handed to a ROCr copy is flushed when its use begins, since ROCr's SDMA and blit rings are in system memory. - The async handler holds a reference to the signal until it has processed it, and drops it before the VirtualGPU reference and inside the async-handler scope, so neither the signal nor the HW queue is destroyed on the handler thread outside the deferral path. With DEBUG_CLR_DEV_SIGNALS only. With the flag off, every queue and ManagedBuffer signal is created with hsa_signal_create and a direct-dispatch queue batches DEBUG_CLR_MAX_BATCH_SIZE commands, as before: - HwQueueTracker::Create builds the completion-signal ring (64) and the interrupt pool (5) as two device-memory batches. - The kernarg and staging ManagedBuffer chunk signals are a device-memory batch too. A chunk signal there is re-armed with a direct store, flushed only when the AQL ring is in system memory, since a store through ROCr reads it back across PCIe. - A direct-dispatch queue completes its submission batch after 64 commands instead of DEBUG_CLR_MAX_BATCH_SIZE (1000), unless that flag is set (Device::SubmissionBatchLimit). A hipEventRecord marker doesn't own a batch, so in a record/sync loop every command keeps its signal referenced until the batch completes. The tracker parks those signals, and refills its pools with device allocations inside the timed intervals. Shared changes. These apply with DEBUG_CLR_DEV_SIGNALS off too, so turning the flag off does not restore the previous behavior in these paths: - ActiveSignal takes the AQL slot and barrier bit of the packet the signal completes; dispatchBarrierValuePacket reserves its slot before taking the signal. - Timing of a use known to be complete (after a CPU wait, or proven retired) skips the signal value check in CacheTimingData. - A ROCr submission that fails after issuing some of the sub-operations sharing one signal (DMA copyBufferRect line by line, SVM prefetch batch) waits for the issued ones (WaitForIssuedOps) before resetting the slot, with either tracker: their late decrements would otherwise complete the slot's next use early. - WaitForSignal no longer loads the signal before waiting, since the wait loads it first, and a blocking wait returns as soon as its 100 us active phase sees completion instead of also running the blocked wait. - The small-batch sync path (#13532) reads the marker's signal once: CompleteCpuWait checks the GPU-error state directly instead of reloading the signal, and HostQueue::finish skips the HW-event wait for a marker that enqueue already completed. Test (MI300X): - HIP stream/graph test with an LD_PRELOAD batch counter passes with DEBUG_CLR_DEV_SIGNALS=0 and 1; graph_bench --verify passes. - With DEBUG_CLR_DEV_SIGNALS=1, median hipEventSynchronize after an empty kernel drops from 7.2 to 4.7 us (10.9 to 4.7 us for blocking-sync events), and hipStreamSynchronize on completed work from 7.5 to 6.1 us; host-memory signals are unchanged. - Unit_hipStreamSynchronize_ConcurrentSubmitFromOtherThread (#13532) passes with the flag off and on, with DEBUG_CLR_BATCH_CPU_SYNC_SIZE at its default and at 0. - hip_launch_latency -k -g, plain and with GPU_CLR_PROFILE_OUTPUT (a signal per dispatch), completes with --verify counts correct, with the flag off and on. With GPU_CLR_PROFILE_OUTPUT and the flag on, graph enqueue takes 32-42 us (44-47 us with host signals) and a single launch measured with events 5.3-5.5 us (6.9 us). - A hipEventRecord/kernel/hipEventRecord/hipEventSynchronize loop with the flag on allocates one spare batch over 3000 iterations instead of 11. Its mean event time is 4.2 us (5.6-5.8 us with host signals), and 6.2-6.4 us under GPU_CLR_PROFILE_OUTPUT (7.3-7.5 us), with 1-3 of 2000 iterations above 50 us instead of 10-25. - Confirming the fence costs nothing measurable: hip_launch_latency -k -g and a 12-stream test sharing 4 HW queues (6.3-6.4 us per launch) match the read-index-only version. On MI300X the read index never passed an incomplete packet with a spinning kernel ahead of a barrier-AND, barrier-value or kernel fence. - hip-tests EventTest, StreamTest, GraphsTest1, GraphsTest2: failure sets identical to develop with the flag off and on (all pre-existing: order-dependent sticky errors, and a stream capture a graph test leaves open). The graph memory tests that failed on Windows (Unit_hipGraphAddMemAllocNode_Negative_With_Cloneed_Graph, Unit_hipGraphMem_Alloc_Free_NodeGetParams_Negative, Unit_hipGraphAllocNodeMemReuse_MemSteal_Remap) pass with the flag off and on. - With DEBUG_CLR_DEV_SIGNALS=0 no signal goes through hsa_amd_signal_batch_create (checked with an LD_PRELOAD counter). - With every device-memory batch failing (an LD_PRELOAD shim), a 3000-iteration event loop allocates one spare batch and keeps it; before, it allocated 98 batches and grew the host pool to about 6000 signals. - rocvirtual.cpp, rocdevice.cpp and rocblit.cpp compile with ROCR_DYN_DLL defined. JIRA ID : AIRUNTIME-2943
saleelk
added a commit
that referenced
this pull request
Oct 29, 2026
…(part 2) Part 2 of a series moving CLR completion signals to the ROCr batch API and device memory. This part covers the queue completion signals. DEBUG_CLR_DEV_SIGNALS (default off) places each queue's completion signals in device memory when the device has a large BAR and no profiler tool is attached, the same condition as the graph signals, in the same pool (uncached on gfx94x/gfx95x, coarse-grained on gfx125x). Every host access to those signals then crosses PCIe, so VirtualGPU uses a new RetiringSignalTracker, which proves that a signal's last use finished from host-side state instead of reading the signal, and re-arms signals ahead of use in batches. Otherwise VirtualGPU keeps LegacySignalTracker. Device-memory signals, graph and queue, are HIP-only: OpenCL keeps host signals and LegacySignalTracker. RetiringSignalTracker: - Retirement. Each use records its AQL slot, whether the packet had the barrier bit, and the HW queue generation. A use published with the barrier bit is a fence: its completion proves that every earlier packet on the HW queue completed. The packet processor may advance the read index before a packet completes (HSA System Architecture 2.8.3), so the read index only nominates the newest fence it passed, and one load of that fence's signal confirms it. Reading the newest fence in a batch (ProbeFence) likewise retires the whole batch with one load. A HW queue change (NoteQueueChange) starts a new generation. - Arming. Reclaim walks the ring from the oldest use, re-arms retired slots with direct stores and one HDP flush per 32 slots, and stops at the first busy one. A slot whose signal is still shared (an event, a GPU wait or a pending handler) gets an armed spare instead; the old signal is parked and reused once the tracker holds the last reference. A ring about to reuse a busy slot grows, up to 16 x ROC_SIGNAL_POOL_SIZE, before waiting on the oldest signal. - Timing. A use's timing is collected before its signal is re-armed, so no Timestamp lists an armed signal. A command whose signals wrap the ring, such as a graph launch under the built-in profiler, would otherwise have its async handler wait on a re-armed signal while holding the Timestamp the next use needs, and read the next use's timestamps. - Timing-only uses. A dispatch that takes a signal only because the built-in profiler times it gets an armed host-memory signal from a separate pool (PlaceForUse), so the handler reads its value and timestamps without crossing PCIe. A use that a GPU waiter or a sync may read swaps a device-memory signal back into the slot. A slot keeps its signal when the replacement's allocation fell back to host memory. - GPU waits. Signals returned by WaitingSignal for a GPU waiter are pinned with a reference until the consuming packets have passed and every earlier use has retired, so re-arming cannot hide the old use from the waiter. A drained queue (NoteQueueDrained) releases them. - HDP flushes. When the device's AQL rings are in device memory, the arming stores are not flushed: the packets that use the signals reach the ring through the same BAR after the stores, in order, and a waiter on another device loads the value before it waits. A signal handed to a ROCr copy is flushed when its use begins, since ROCr's SDMA and blit rings are in system memory. - The async handler holds a reference to the signal until it has processed it, and drops it before the VirtualGPU reference and inside the async-handler scope, so neither the signal nor the HW queue is destroyed on the handler thread outside the deferral path. With DEBUG_CLR_DEV_SIGNALS only. With the flag off, every queue and ManagedBuffer signal is created with hsa_signal_create and a direct-dispatch queue batches DEBUG_CLR_MAX_BATCH_SIZE commands, as before: - HwQueueTracker::Create builds the completion-signal ring (64) and the interrupt pool (5) as two device-memory batches. - The kernarg and staging ManagedBuffer chunk signals are a device-memory batch too. A chunk signal there is re-armed with a direct store, flushed only when the AQL ring is in system memory, since a store through ROCr reads it back across PCIe. - A direct-dispatch queue completes its submission batch after 64 commands instead of DEBUG_CLR_MAX_BATCH_SIZE (1000), unless that flag is set (Device::SubmissionBatchLimit). A hipEventRecord marker doesn't own a batch, so in a record/sync loop every command keeps its signal referenced until the batch completes. The tracker parks those signals, and refills its pools with device allocations inside the timed intervals. Shared changes. These apply with DEBUG_CLR_DEV_SIGNALS off too, so turning the flag off does not restore the previous behavior in these paths: - ActiveSignal takes the AQL slot and barrier bit of the packet the signal completes; dispatchBarrierValuePacket reserves its slot before taking the signal. - Timing of a use known to be complete (after a CPU wait, or proven retired) skips the signal value check in CacheTimingData. - A ROCr submission that fails after issuing some of the sub-operations sharing one signal (DMA copyBufferRect line by line, SVM prefetch batch) waits for the issued ones (WaitForIssuedOps) before resetting the slot, with either tracker: their late decrements would otherwise complete the slot's next use early. - WaitForSignal no longer loads the signal before waiting, since the wait loads it first, and a blocking wait returns as soon as its 100 us active phase sees completion instead of also running the blocked wait. - The small-batch sync path (#13532) reads the marker's signal once: CompleteCpuWait checks the GPU-error state directly instead of reloading the signal, and HostQueue::finish skips the HW-event wait for a marker that enqueue already completed. Test (MI300X, container_saleel_1, image rocm/ufb-private:vllm-0.26.0-ubuntu24.04-py3.14-nightlies-device-all-cdna-rocm10.1.0a20260826-568afb3a13, ROCm SDK 10.1.0a20260826, ROCr built from this PR's develop base). Host time per call in us, median of 2000: | API | develop | DEBUG_CLR_DEV_SIGNALS=0 | DEBUG_CLR_DEV_SIGNALS=1 | |---|---|---|---| | hipEventSynchronize (empty kernel, record, sync) | 5.3-5.6 | 5.3-5.6 | 4.8-5.0 | | hipEventQuery (completed event) | 0.08 | 0.08 | 1.2-1.3 | | hipStreamSynchronize (empty kernel, sync) | 6.5-6.9 | 6.5-6.9 | 5.6-5.8 | | hipDeviceSynchronize (empty kernel, sync) | 6.5-6.9 | 6.5-6.9 | 5.7-5.8 | - hip_launch_latency -k -g --verify passes with the flag off and on, with and without GPU_CLR_PROFILE_OUTPUT. - hip-tests EventTest, StreamTest, GraphsTest1, GraphsTest2: failure sets identical to develop with the flag off and on. - rocvirtual.cpp, rocdevice.cpp and rocblit.cpp compile with ROCR_DYN_DLL defined. JIRA ID : AIRUNTIME-2943
saleelk
added a commit
that referenced
this pull request
Oct 29, 2026
…(part 2) Part 2 of a series moving CLR completion signals to the ROCr batch API and device memory. This part covers the queue completion signals. DEBUG_CLR_DEV_SIGNALS (default off) places each queue's completion signals in device memory when the device has a large BAR and no profiler tool is attached, the same condition as the graph signals, in the same pool (uncached on gfx94x/gfx95x, coarse-grained on gfx125x). Every host access to those signals then crosses PCIe, so VirtualGPU uses a new RetiringSignalTracker, which proves that a signal's last use finished from host-side state instead of reading the signal, and re-arms signals ahead of use in batches. Otherwise VirtualGPU keeps LegacySignalTracker. Device-memory signals, graph and queue, are HIP-only: OpenCL keeps host signals and LegacySignalTracker. RetiringSignalTracker: - Retirement. Each use records its AQL slot, whether the packet had the barrier bit, and the HW queue generation. A use published with the barrier bit is a fence: its completion proves that every earlier packet on the HW queue completed. The packet processor may advance the read index before a packet completes (HSA System Architecture 2.8.3), so the read index only nominates the newest fence it passed, and one load of that fence's signal confirms it. Reading the newest fence in a batch (ProbeFence) likewise retires the whole batch with one load. A HW queue change (NoteQueueChange) starts a new generation. - Arming. Reclaim walks the ring from the oldest use, re-arms retired slots with direct stores and one HDP flush per 32 slots, and stops at the first busy one. A slot whose signal is still shared (an event, a GPU wait or a pending handler) gets an armed spare instead; the old signal is parked and reused once the tracker holds the last reference. A ring about to reuse a busy slot grows, up to 16 x ROC_SIGNAL_POOL_SIZE, before waiting on the oldest signal. - Timing. A use's timing is collected before its signal is re-armed, so no Timestamp lists an armed signal. A command whose signals wrap the ring, such as a graph launch under the built-in profiler, would otherwise have its async handler wait on a re-armed signal while holding the Timestamp the next use needs, and read the next use's timestamps. - Timing-only uses. A dispatch that takes a signal only because the built-in profiler times it gets an armed host-memory signal from a separate pool (PlaceForUse), so the handler reads its value and timestamps without crossing PCIe. A use that a GPU waiter or a sync may read swaps a device-memory signal back into the slot. A slot keeps its signal when the replacement's allocation fell back to host memory. - GPU waits. Signals returned by WaitingSignal for a GPU waiter are pinned with a reference until the consuming packets have passed and every earlier use has retired, so re-arming cannot hide the old use from the waiter. A drained queue (NoteQueueDrained) releases them. - Store ordering. Host stores to device-memory signals are write-combined. FlushStores fences them, so a peer GPU that waits on a signal over XGMI reads the re-arm, and flushes HDP only when the AQL rings are in system memory: a device-memory ring receives the packets that use the signals through the same BAR after the stores, in order. A signal handed to a ROCr copy is flushed when its use begins, since ROCr's SDMA and blit rings are in system memory. - The async handler holds a reference to the signal until it has processed it, and drops it before the VirtualGPU reference and inside the async-handler scope, so neither the signal nor the HW queue is destroyed on the handler thread outside the deferral path. With DEBUG_CLR_DEV_SIGNALS only. With the flag off, every queue and ManagedBuffer signal is created with hsa_signal_create and a direct-dispatch queue batches DEBUG_CLR_MAX_BATCH_SIZE commands, as before: - HwQueueTracker::Create builds the completion-signal ring (64) and the interrupt pool (5) as two device-memory batches. - The kernarg and staging ManagedBuffer chunk signals are a device-memory batch too. A chunk signal there is re-armed with a direct store, flushed only when the AQL ring is in system memory, since a store through ROCr reads it back across PCIe. - A direct-dispatch queue completes its submission batch after 64 commands instead of DEBUG_CLR_MAX_BATCH_SIZE (1000), unless that flag is set (Device::SubmissionBatchLimit). A hipEventRecord marker doesn't own a batch, so in a record/sync loop every command keeps its signal referenced until the batch completes. The tracker parks those signals, and refills its pools with device allocations inside the timed intervals. - WaitForSignal skips its check before waiting for a device-memory signal, since the wait loads the value first and each load is a PCIe read. Host-memory signals keep the check. Shared changes. These apply with DEBUG_CLR_DEV_SIGNALS off too: - ActiveSignal takes the AQL slot and barrier bit of the packet the signal completes; dispatchBarrierValuePacket reserves its slot before taking the signal. - Timing of a use known to be complete (after a CPU wait, or proven retired) skips the signal value check in CacheTimingData. - A ROCr submission that fails after issuing some of the sub-operations sharing one signal (DMA copyBufferRect line by line, SVM prefetch batch) waits for the issued ones (WaitForIssuedOps) before resetting the slot, with either tracker: their late decrements would otherwise complete the slot's next use early. - A blocking WaitForSignal returns as soon as its 100 us active phase sees completion, instead of also running the blocked wait. - The small-batch sync path (#13532) reads the marker's signal once: CompleteCpuWait checks the GPU-error state directly instead of reloading the signal, and HostQueue::finish skips the HW-event wait for a marker that enqueue already completed. Known limitations with DEBUG_CLR_DEV_SIGNALS=1: - The event APIs read the event's signal and timestamps across PCIe, about 1.1 us per read (EventPerformance table below). The gain is in waits: completion lands in device memory, and GPU waiters poll local memory. - hipEventBlockingSync events spin, as they do with host signals: ROCr's busy-wait signals ignore the wait hint. With the flag on, each poll is a PCIe read. - Validated on MI300X only (PCIe to the host, XGMI between GPUs). IPC events are unaffected: they complete through a separate ROCr IPC signal. Test (MI300X, container_saleel_1, image rocm/ufb-private:vllm-0.26.0-ubuntu24.04-py3.14-nightlies-device-all-cdna-rocm10.1.0a20260826-568afb3a13, ROCm SDK 10.1.0a20260826, ROCr built from this PR's develop base). Host time per call in us, median of 2000: | API | develop | DEBUG_CLR_DEV_SIGNALS=0 | DEBUG_CLR_DEV_SIGNALS=1 | |---|---|---|---| | hipEventSynchronize (empty kernel, record, sync) | 5.3 | 5.3 | 4.5-4.7 | | hipEventSynchronize, blocking-sync event | 5.4 | 5.3 | 4.6-4.7 | | hipEventQuery (completed event) | 0.08 | 0.08 | 1.2 | | hipStreamSynchronize (empty kernel, sync) | 6.5 | 6.5 | 6.3-6.4 | | hipStreamSynchronize (completed work) | 0.14 | 0.14 | 0.14 | | hipDeviceSynchronize (empty kernel, sync) | 6.5 | 6.5 | 6.3-6.5 | hip-tests EventPerformance (HIP_TEST_LEVEL=level_1), average in us. Each timed call includes the harness's hipStreamSynchronize: | Benchmark | develop | DEBUG_CLR_DEV_SIGNALS=0 | DEBUG_CLR_DEV_SIGNALS=1 | |---|---|---|---| | Performance_hipEventRecord | 6.6 | 6.6 | 8.2-8.5 | | Performance_hipEventSynchronize, default | 6.1-6.4 | 6.3 | 8.5-8.9 | | Performance_hipEventSynchronize, blocking | 6.2 | 6.1 | 8.5-8.8 | | Performance_hipEventElapsedTime | 1.4 | 1.4 | 16.6-17.2 | | Performance_hipEventQuery | 0.78 | 0.78 | 5.0-5.2 | - hip_launch_latency -k -g --verify passes with the flag off and on, with and without GPU_CLR_PROFILE_OUTPUT. - hipStreamWaitEvent on GPU1 for a GPU0 event recorded behind a spinning kernel never passed early (0 of 300 rounds), with the flag off and on. - hip-tests EventTest, StreamTest, GraphsTest1, GraphsTest2: failure sets identical to develop with the flag off and on. - rocvirtual.cpp, rocdevice.cpp and rocblit.cpp compile with ROCR_DYN_DLL defined. JIRA ID : AIRUNTIME-2943
renjithravindrankannath
pushed a commit
that referenced
this pull request
Oct 29, 2026
…(part 2) (#13507) Part 2 of a series moving CLR completion signals to the ROCr batch API and device memory. This part covers the queue completion signals. DEBUG_CLR_DEV_SIGNALS (default off) places each queue's completion signals in device memory when the device has a large BAR and no profiler tool is attached, the same condition as the graph signals, in the same pool (uncached on gfx94x/gfx95x, coarse-grained on gfx125x). Every host access to those signals then crosses PCIe, so VirtualGPU uses a new RetiringSignalTracker, which proves that a signal's last use finished from host-side state instead of reading the signal, and re-arms signals ahead of use in batches. Otherwise VirtualGPU keeps LegacySignalTracker. Device-memory signals, graph and queue, are HIP-only: OpenCL keeps host signals and LegacySignalTracker. RetiringSignalTracker: - Retirement. Each use records its AQL slot, whether the packet had the barrier bit, and the HW queue generation. A use published with the barrier bit is a fence: its completion proves that every earlier packet on the HW queue completed. The packet processor may advance the read index before a packet completes (HSA System Architecture 2.8.3), so the read index only nominates the newest fence it passed, and one load of that fence's signal confirms it. Reading the newest fence in a batch (ProbeFence) likewise retires the whole batch with one load. A HW queue change (NoteQueueChange) starts a new generation. - Arming. Reclaim walks the ring from the oldest use, re-arms retired slots with direct stores and one HDP flush per 32 slots, and stops at the first busy one. A slot whose signal is still shared (an event, a GPU wait or a pending handler) gets an armed spare instead; the old signal is parked and reused once the tracker holds the last reference. A ring about to reuse a busy slot grows, up to 16 x ROC_SIGNAL_POOL_SIZE, before waiting on the oldest signal. - Timing. A use's timing is collected before its signal is re-armed, so no Timestamp lists an armed signal. A command whose signals wrap the ring, such as a graph launch under the built-in profiler, would otherwise have its async handler wait on a re-armed signal while holding the Timestamp the next use needs, and read the next use's timestamps. - Timing-only uses. A dispatch that takes a signal only because the built-in profiler times it gets an armed host-memory signal from a separate pool (PlaceForUse), so the handler reads its value and timestamps without crossing PCIe. A use that a GPU waiter or a sync may read swaps a device-memory signal back into the slot. A slot keeps its signal when the replacement's allocation fell back to host memory. - GPU waits. Signals returned by WaitingSignal for a GPU waiter are pinned with a reference until the consuming packets have passed and every earlier use has retired, so re-arming cannot hide the old use from the waiter. A drained queue (NoteQueueDrained) releases them. - Store ordering. Host stores to device-memory signals are write-combined. FlushStores fences them, so a peer GPU that waits on a signal over XGMI reads the re-arm, and flushes HDP only when the AQL rings are in system memory: a device-memory ring receives the packets that use the signals through the same BAR after the stores, in order. A signal handed to a ROCr copy is flushed when its use begins, since ROCr's SDMA and blit rings are in system memory. - The async handler holds a reference to the signal until it has processed it, and drops it before the VirtualGPU reference and inside the async-handler scope, so neither the signal nor the HW queue is destroyed on the handler thread outside the deferral path. With DEBUG_CLR_DEV_SIGNALS only. With the flag off, every queue and ManagedBuffer signal is created with hsa_signal_create and a direct-dispatch queue batches DEBUG_CLR_MAX_BATCH_SIZE commands, as before: - HwQueueTracker::Create builds the completion-signal ring (64) and the interrupt pool (5) as two device-memory batches. - The kernarg and staging ManagedBuffer chunk signals are a device-memory batch too. A chunk signal there is re-armed with a direct store, flushed only when the AQL ring is in system memory, since a store through ROCr reads it back across PCIe. - A direct-dispatch queue completes its submission batch after 64 commands instead of DEBUG_CLR_MAX_BATCH_SIZE (1000), unless that flag is set (Device::SubmissionBatchLimit). A hipEventRecord marker doesn't own a batch, so in a record/sync loop every command keeps its signal referenced until the batch completes. The tracker parks those signals, and refills its pools with device allocations inside the timed intervals. - WaitForSignal skips its check before waiting for a device-memory signal, since the wait loads the value first and each load is a PCIe read. Host-memory signals keep the check. Shared changes. These apply with DEBUG_CLR_DEV_SIGNALS off too: - ActiveSignal takes the AQL slot and barrier bit of the packet the signal completes; dispatchBarrierValuePacket reserves its slot before taking the signal. - Timing of a use known to be complete (after a CPU wait, or proven retired) skips the signal value check in CacheTimingData. - A ROCr submission that fails after issuing some of the sub-operations sharing one signal (DMA copyBufferRect line by line, SVM prefetch batch) waits for the issued ones (WaitForIssuedOps) before resetting the slot, with either tracker: their late decrements would otherwise complete the slot's next use early. - A blocking WaitForSignal returns as soon as its 100 us active phase sees completion, instead of also running the blocked wait. - The small-batch sync path (#13532) reads the marker's signal once: CompleteCpuWait checks the GPU-error state directly instead of reloading the signal, and HostQueue::finish skips the HW-event wait for a marker that enqueue already completed. Known limitations with DEBUG_CLR_DEV_SIGNALS=1: - The event APIs read the event's signal and timestamps across PCIe, about 1.1 us per read (EventPerformance table below). The gain is in waits: completion lands in device memory, and GPU waiters poll local memory. - hipEventBlockingSync events spin, as they do with host signals: ROCr's busy-wait signals ignore the wait hint. With the flag on, each poll is a PCIe read. - Validated on MI300X only (PCIe to the host, XGMI between GPUs). IPC events are unaffected: they complete through a separate ROCr IPC signal. Test (MI300X, container_saleel_1, image rocm/ufb-private:vllm-0.26.0-ubuntu24.04-py3.14-nightlies-device-all-cdna-rocm10.1.0a20260826-568afb3a13, ROCm SDK 10.1.0a20260826, ROCr built from this PR's develop base). Host time per call in us, median of 2000: | API | develop | DEBUG_CLR_DEV_SIGNALS=0 | DEBUG_CLR_DEV_SIGNALS=1 | |---|---|---|---| | hipEventSynchronize (empty kernel, record, sync) | 5.3 | 5.3 | 4.5-4.7 | | hipEventSynchronize, blocking-sync event | 5.4 | 5.3 | 4.6-4.7 | | hipEventQuery (completed event) | 0.08 | 0.08 | 1.2 | | hipStreamSynchronize (empty kernel, sync) | 6.5 | 6.5 | 6.3-6.4 | | hipStreamSynchronize (completed work) | 0.14 | 0.14 | 0.14 | | hipDeviceSynchronize (empty kernel, sync) | 6.5 | 6.5 | 6.3-6.5 | hip-tests EventPerformance (HIP_TEST_LEVEL=level_1), average in us. Each timed call includes the harness's hipStreamSynchronize: | Benchmark | develop | DEBUG_CLR_DEV_SIGNALS=0 | DEBUG_CLR_DEV_SIGNALS=1 | |---|---|---|---| | Performance_hipEventRecord | 6.6 | 6.6 | 8.2-8.5 | | Performance_hipEventSynchronize, default | 6.1-6.4 | 6.3 | 8.5-8.9 | | Performance_hipEventSynchronize, blocking | 6.2 | 6.1 | 8.5-8.8 | | Performance_hipEventElapsedTime | 1.4 | 1.4 | 16.6-17.2 | | Performance_hipEventQuery | 0.78 | 0.78 | 5.0-5.2 | - hip_launch_latency -k -g --verify passes with the flag off and on, with and without GPU_CLR_PROFILE_OUTPUT. - hipStreamWaitEvent on GPU1 for a GPU0 event recorded behind a spinning kernel never passed early (0 of 300 rounds), with the flag off and on. - hip-tests EventTest, StreamTest, GraphsTest1, GraphsTest2: failure sets identical to develop with the flag off and on. - rocvirtual.cpp, rocdevice.cpp and rocblit.cpp compile with ROCR_DYN_DLL defined. JIRA ID : AIRUNTIME-2943
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
hipStreamSynchronize() can block every other thread that uses the same stream, and a process whose GPU work depends on one of those threads might deadlock (for example torch.distributed gloo block_current_stream).
HostQueue::finish() sends a Marker with cpu_wait=true when the pending batch is smaller than DEBUG_CLR_BATCH_CPU_SYNC_SIZE (the small-batch path from #5176). VirtualGPU::submitMarker() flushed and busy-waited on the HSA signal inside Command::enqueue(), i.e. while holding the queue's execution() lock. Until the GPU finished, every other thread enqueuing to the same stream was blocked: hipStreamWaitEvent, kernel launches, hipMemcpyAsync, and hipStreamQuery (through getLastQueuedCommand).
Technical Details
The lock only has to cover the submission, not the GPU execution. Command::enqueue() keeps its scoped_lock and the batch/ring ordering it protects; only the wait moves out of it.
Issue Tracking
GitHub issue: #11678
JIRA ID: ROCM-32517
Test Plan
Added Unit_hipStreamSynchronize_ConcurrentSubmitFromOtherThread, with a device-side watchdog so a deadlock becomes a test failure. On gfx1201 it fails without the fix and passes with it. The reproducer from the issue shows no async call blocked for more than 50 ms. Multi-thread sync stress shows no early return or deadlock. Single-thread finish() latency is unchanged. StreamTest/EventTest have no test that fails only with the fix.
Test Result
All passed
Submission Checklist