fix(clr): Order shared AQL ring publication on Intel hosts (AIRUNTIME-2291) - #282
Open
saleelk wants to merge 1 commit into
Open
fix(clr): Order shared AQL ring publication on Intel hosts (AIRUNTIME-2291)#282saleelk wants to merge 1 commit into
saleelk wants to merge 1 commit into
Conversation
saleelk
requested review from
chrispaquot,
mangupta and
rakesroy
as code owners
August 18, 2026 23:38
saleelk
force-pushed
the
users/saleelk/orderedDB_63
branch
2 times, most recently
from
August 19, 2026 15:55
17d44f3 to
91f474e
Compare
When several VirtualGPUs multiplex onto one physical AQL ring, hsa_queue_add_write_index_screlease() reserves a slot and advances the GPU-visible write_dispatch_id in one step, so the ring advertises slots whose packets are still being written. CPF is sent to fetch such a slot by a doorbell from a producer holding a later slot, and a 64B packet fetch that splits into smaller reads can assemble a valid header over a stale body, so MEC runs the packet with a junk body. Order the doorbell rather than the reservation: slots are still claimed with one atomic add, so producers own disjoint slots and fill them in parallel, but a producer holds its doorbell until every earlier slot has been rung. Each pooled HW queue carries that doorbell watermark, and it only moves forward, so a write index rewound from outside CLR cannot leave a reservation waiting behind a slot that is gone. CLR never stores the write index, which is what let rocprofiler's virtualized index hang counter collection. DEBUG_CLR_ORDER_DOORBELL selects which hosts order their doorbells: DEBUG_CLR_ORDER_DOORBELL=0 never orders, =1 orders Intel hosts only, which is the only place the torn fetch has been observed, and =2 orders every host. It defaults to 2, since nothing rules the race out elsewhere. CU masked and cooperative queues are never shared between VirtualGPUs, so they keep the original HSA multi-producer path.
saleelk
force-pushed
the
users/saleelk/orderedDB_63
branch
from
September 2, 2026 05:34
91f474e to
25510bf
Compare
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.
Summary
Backport of the ordered AQL ring publication fix to
release/rocm-rel-6.3.0.1.When several VirtualGPUs multiplex onto one physical AQL ring,
hsa_queue_add_write_index_screleaseadvances the GPU-visible write index at reservation time. A later producer can therefore expose
its slot while an earlier producer is still writing its packet, and the in-order command processor
parks on that unwritten packet (or consumes torn state).
This splits slot reservation from publication:
AqlOrderedPublishStatewith shared reserve/commit indices.VirtualGPU::ReserveAqlSlots()allocates slots without making them visible to the GPU.VirtualGPU::CommitAqlSlots()stores the write index and rings the doorbell only once everyolder reservation has been published.
Gated to GenuineIntel hosts and controlled by the new
DEBUG_CLR_ORDER_DOORBELLflag (default on);other hosts keep the original HSA multi-producer path. Cooperative and CU-masked queues are not
shared between VirtualGPUs, so they keep the original path as well.
Publication order equals reservation order, so this cannot deadlock: the only blocking operations a
stream performs while holding a reservation (queue slot waits, completion signal waits) wait on
strictly older packets, which are published by definition.
Notes on this backport
The equivalent change is already on
developandrelease/rocm-rel-7.2. Differences here followfrom the older code base:
rocrctxwrapper layer, so HSA entry points are called directly.Device::QueueInfo(there is noQueueExtras), and is looked up undervgpusAccess_, which is the lock that guards the queue pools on this branch.gpu_queue_is bound exactly once inVirtualGPU::create(), so noset_gpu_queue()helper isneeded to refresh the cached state.
pacing on 7.2 is omitted.
Test plan
-fsyntax-only).all
GPU_MAX_HW_QUEUES/DEBUG_CLR_ORDER_DOORBELLcombinations).