[GDA/IONIC] only ring doorbell on active lanes (#2727)

Co-authored-by: Yiltan <yiltan@amd.com>
Co-authored-by: systems-assistant[bot] <systems-assistant[bot]@users.noreply.github.com>
Co-authored-by: JeniferC99 <150404595+JeniferC99@users.noreply.github.com>
Este commit está contenido en:
systems-assistant[bot]
2026-01-23 09:26:46 -05:00
cometido por GitHub
padre 48347bc857
commit 563776a949
@@ -204,8 +204,12 @@ __device__ void QueuePair::ionic_quiet_internal(uint64_t activemask, uint32_t co
__device__ void QueuePair::ionic_ring_doorbell(uint32_t pos) {
// When threads write at once to the same address, not all writes reach the bus.
// Take turns and insert a thread fence between writes to the same address.
for (int i = 0; i < WF_SIZE; ++i) {
if (__lane_id() == i) {
uint32_t activemask = __activemask();
uint32_t lane_id = get_active_lane_num(activemask);
uint32_t lane_count = get_active_lane_count(activemask);
for (int i = 0; i < lane_count; ++i) {
if (lane_id == i) {
__threadfence();
__atomic_store_n(sq_dbreg, sq_dbval | (sq_mask & pos), __ATOMIC_SEQ_CST);
}