Skip to content

fix(clr): Order shared AQL ring publication on Intel hosts (AIRUNTIME-2291) - #282

Open
saleelk wants to merge 1 commit into
ROCm:release/rocm-rel-6.3.0.1from
saleelk:users/saleelk/orderedDB_63
Open

fix(clr): Order shared AQL ring publication on Intel hosts (AIRUNTIME-2291)#282
saleelk wants to merge 1 commit into
ROCm:release/rocm-rel-6.3.0.1from
saleelk:users/saleelk/orderedDB_63

Conversation

@saleelk

@saleelk saleelk commented Aug 18, 2026

Copy link
Copy Markdown
Contributor

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_screlease
advances 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:

  • Each pooled hardware queue owns an AqlOrderedPublishState with 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 every
    older reservation has been published.

Gated to GenuineIntel hosts and controlled by the new DEBUG_CLR_ORDER_DOORBELL flag (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 develop and release/rocm-rel-7.2. Differences here follow
from the older code base:

  • No rocrctx wrapper layer, so HSA entry points are called directly.
  • The state lives in Device::QueueInfo (there is no QueueExtras), and is looked up under
    vgpusAccess_, which is the lock that guards the queue pools on this branch.
  • gpu_queue_ is bound exactly once in VirtualGPU::create(), so no set_gpu_queue() helper is
    needed to refresh the cached state.
  • There is no batch dispatch path on this branch, so the reservation-index accessor used for batch
    pacing on 7.2 is omitted.

Test plan

  • Each modified translation unit compiles clean (-fsyntax-only).
  • Runtime validation on a 6.3.0.1 stack (the equivalent 7.2 change passed MultiThreadTest across
    all GPU_MAX_HW_QUEUES / DEBUG_CLR_ORDER_DOORBELL combinations).

@saleelk
saleelk force-pushed the users/saleelk/orderedDB_63 branch 2 times, most recently from 17d44f3 to 91f474e Compare August 19, 2026 15:55
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
saleelk force-pushed the users/saleelk/orderedDB_63 branch from 91f474e to 25510bf Compare September 2, 2026 05:34
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant