From 563776a949047ac3846e431cfcb31d707cfd4021 Mon Sep 17 00:00:00 2001 From: "systems-assistant[bot]" <221163467+systems-assistant[bot]@users.noreply.github.com> Date: Fri, 23 Jan 2026 09:26:46 -0500 Subject: [PATCH] [GDA/IONIC] only ring doorbell on active lanes (#2727) Co-authored-by: Yiltan Co-authored-by: systems-assistant[bot] Co-authored-by: JeniferC99 <150404595+JeniferC99@users.noreply.github.com> --- projects/rocshmem/src/gda/ionic/queue_pair_ionic.cpp | 8 ++++++-- 1 file changed, 6 insertions(+), 2 deletions(-) diff --git a/projects/rocshmem/src/gda/ionic/queue_pair_ionic.cpp b/projects/rocshmem/src/gda/ionic/queue_pair_ionic.cpp index fd8a960380..30be0c4a39 100644 --- a/projects/rocshmem/src/gda/ionic/queue_pair_ionic.cpp +++ b/projects/rocshmem/src/gda/ionic/queue_pair_ionic.cpp @@ -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); }