From 93493e3e465e4cd5df12d954b5dfa9f1a90f0744 Mon Sep 17 00:00:00 2001 From: Omri Mor Date: Thu, 15 Jan 2026 15:19:19 -0700 Subject: [PATCH] ionic: fix byteswap functions (added in #345), missed in #368 (#388) [ROCm/rocshmem commit: 885e41ec62151f3eefd5c2391b6c8f39ac6d2d2c] --- projects/rocshmem/src/gda/ionic/queue_pair_ionic.cpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/projects/rocshmem/src/gda/ionic/queue_pair_ionic.cpp b/projects/rocshmem/src/gda/ionic/queue_pair_ionic.cpp index 9a9cea5b7c..164219c410 100644 --- a/projects/rocshmem/src/gda/ionic/queue_pair_ionic.cpp +++ b/projects/rocshmem/src/gda/ionic/queue_pair_ionic.cpp @@ -75,15 +75,15 @@ __device__ void QueuePair::ionic_poll_wave_ccqe(uint64_t activemask) { struct ionic_v1_cqe *cqe = &ionic_cq_buf[0]; uint32_t qtf_be = *(volatile uint32_t *)(&cqe->qid_type_flags); - uint32_t msn = swap_endian_val(cqe->send.msg_msn); + uint32_t msn = byteswap(cqe->send.msg_msn); - if (!!(qtf_be & swap_endian_val(IONIC_V1_CQE_ERROR))) { + if (!!(qtf_be & byteswap(IONIC_V1_CQE_ERROR))) { #if defined(DEBUG) - uint32_t qtf = swap_endian_val(qtf_be); + uint32_t qtf = byteswap(qtf_be); uint32_t qid = qtf >> IONIC_V1_CQE_QID_SHIFT; uint32_t type = (qtf >> IONIC_V1_CQE_TYPE_SHIFT) & IONIC_V1_CQE_TYPE_MASK; uint32_t flag = qtf & 0xf; - uint32_t status = swap_endian_val(cqe->status_length); + uint32_t status = byteswap(cqe->status_length); uint64_t npg = cqe->send.npg_wqe_idx_timestamp & IONIC_V1_CQE_WQE_IDX_MASK; printf("QUIET ERROR (CCQE): %s qid %u type %u flag %#x status %u msn %u npg %lu\n", @@ -169,9 +169,9 @@ __device__ void QueuePair::ionic_quiet_internal_ccqe(uint64_t activemask, uint32 } volatile struct ionic_v1_cqe *cqe = &ionic_cq_buf[0]; - uint32_t msn = swap_endian_val(cqe->send.msg_msn); + uint32_t msn = byteswap(cqe->send.msg_msn); while ((msn - cons) & 0x800000) { - msn = swap_endian_val(cqe->send.msg_msn); + msn = byteswap(cqe->send.msg_msn); } }