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); }