From 71583811b17fcd3008c29a5ef94e5296a2e9228a Mon Sep 17 00:00:00 2001 From: qizzhang Date: Wed, 6 May 2026 11:46:13 +0800 Subject: [PATCH 01/33] replace ionic header files with ccqe version --- .../transport/rdma/providers/ionic/ionic_dv.h | 35 + .../transport/rdma/providers/ionic/ionic_fw.h | 685 ++++++++++-------- 2 files changed, 409 insertions(+), 311 deletions(-) diff --git a/include/mori/core/transport/rdma/providers/ionic/ionic_dv.h b/include/mori/core/transport/rdma/providers/ionic/ionic_dv.h index a1815990..faa167ec 100644 --- a/include/mori/core/transport/rdma/providers/ionic/ionic_dv.h +++ b/include/mori/core/transport/rdma/providers/ionic/ionic_dv.h @@ -177,6 +177,16 @@ int ionic_dv_pd_set_sqcmb(struct ibv_pd* ibpd, bool enable, bool expdb, bool req */ int ionic_dv_pd_set_rqcmb(struct ibv_pd* ibpd, bool enable, bool expdb, bool require); +/** + * ionic_dv_pd_set_expdb_mask - Specify expdb mask. + * + * Queues associated with this pd will attempt to have expdb on for WQE sizes + * other than default (and supported by the NIC). + * + * @mask - IONIC_EPXDB_* bitmap + */ +int ionic_dv_pd_set_expdb_mask(struct ibv_pd* ibpd, uint8_t mask); + /** * ionic_dv_qp_set_gda - Enable or disable GPU-Direct Async (GDA) mode. * @@ -241,6 +251,31 @@ int ionic_dv_qp_get_send_dbell_data(struct ibv_qp* ibqp, uint64_t* dbdata); */ int ionic_dv_qp_get_recv_dbell_data(struct ibv_qp* ibqp, uint64_t* dbdata); +enum ionic_cq_init_attr_mask { + IONIC_CQ_INIT_ATTR_MASK_FLAGS = 1 << 0, +}; + +enum ionic_cq_init_attr_flags { + IONIC_CQ_INIT_ATTR_CCQE = 1 << 0, +}; + +struct ionic_cq_init_attr_ex { + /* One or more flags of enum ionic_cq_init_attr_mask */ + uint32_t comp_mask; + /* One or more flags of enum ionic_cq_init_attr_flags */ + uint32_t flags; +}; + +/** + * ionic_dv_create_cq_ex - Create an IBV CQ with vendor-specific attributes. + * + * @ibctx - Context CQ will be attached to. + * @ex - IBV attributes to create the CQ with. + * @ionic_ex - Vendor-specific attributes to create the CQ with. + */ +struct ibv_cq_ex* ionic_dv_create_cq_ex(struct ibv_context* ibctx, struct ibv_cq_init_attr_ex* ex, + struct ionic_cq_init_attr_ex* ionic_ex); + /** * ionic_dv_get_ctx - Extract context information for gpu-initiated rdma. */ diff --git a/include/mori/core/transport/rdma/providers/ionic/ionic_fw.h b/include/mori/core/transport/rdma/providers/ionic/ionic_fw.h index da5d5371..13000476 100644 --- a/include/mori/core/transport/rdma/providers/ionic/ionic_fw.h +++ b/include/mori/core/transport/rdma/providers/ionic/ionic_fw.h @@ -42,468 +42,531 @@ /* wqe scatter gather element */ struct ionic_sge { - __be64 va; - __be32 len; - __be32 lkey; + __be64 va; + __be32 len; + __be32 lkey; }; /* admin queue mr type */ enum ionic_mr_flags { - /* bits that determine mr access */ - IONIC_MRF_LOCAL_WRITE = BIT(0), - IONIC_MRF_REMOTE_WRITE = BIT(1), - IONIC_MRF_REMOTE_READ = BIT(2), - IONIC_MRF_REMOTE_ATOMIC = BIT(3), - IONIC_MRF_MW_BIND = BIT(4), - IONIC_MRF_ZERO_BASED = BIT(5), - IONIC_MRF_ON_DEMAND = BIT(6), - IONIC_MRF_PB = BIT(7), - IONIC_MRF_ACCESS_MASK = BIT(12) - 1, - - /* bits that determine mr type */ - IONIC_MRF_IS_MW = BIT(14), - IONIC_MRF_INV_EN = BIT(15), - - /* base flags combinations for mr types */ - IONIC_MRF_USER_MR = 0, - IONIC_MRF_PHYS_MR = IONIC_MRF_INV_EN, - IONIC_MRF_MW_1 = IONIC_MRF_IS_MW, - IONIC_MRF_MW_2 = IONIC_MRF_IS_MW | IONIC_MRF_INV_EN, + /* bits that determine mr access */ + IONIC_MRF_LOCAL_WRITE = BIT(0), + IONIC_MRF_REMOTE_WRITE = BIT(1), + IONIC_MRF_REMOTE_READ = BIT(2), + IONIC_MRF_REMOTE_ATOMIC = BIT(3), + IONIC_MRF_MW_BIND = BIT(4), + IONIC_MRF_ZERO_BASED = BIT(5), + IONIC_MRF_ON_DEMAND = BIT(6), + IONIC_MRF_PB = BIT(7), + IONIC_MRF_ACCESS_MASK = BIT(12) - 1, + + /* bits that determine mr type */ + IONIC_MRF_IS_MW = BIT(14), + IONIC_MRF_INV_EN = BIT(15), + + /* base flags combinations for mr types */ + IONIC_MRF_USER_MR = 0, + IONIC_MRF_PHYS_MR = IONIC_MRF_INV_EN, + IONIC_MRF_MW_1 = IONIC_MRF_IS_MW, + IONIC_MRF_MW_2 = IONIC_MRF_IS_MW | IONIC_MRF_INV_EN, }; -static inline int to_ionic_mr_flags(int access) { - int flags = 0; +static inline int to_ionic_mr_flags(int access) +{ + int flags = 0; - if (access & IBV_ACCESS_LOCAL_WRITE) flags |= IONIC_MRF_LOCAL_WRITE; + if (access & IBV_ACCESS_LOCAL_WRITE) + flags |= IONIC_MRF_LOCAL_WRITE; - if (access & IBV_ACCESS_REMOTE_READ) flags |= IONIC_MRF_REMOTE_READ; + if (access & IBV_ACCESS_REMOTE_READ) + flags |= IONIC_MRF_REMOTE_READ; - if (access & IBV_ACCESS_REMOTE_WRITE) flags |= IONIC_MRF_REMOTE_WRITE; + if (access & IBV_ACCESS_REMOTE_WRITE) + flags |= IONIC_MRF_REMOTE_WRITE; - if (access & IBV_ACCESS_REMOTE_ATOMIC) flags |= IONIC_MRF_REMOTE_ATOMIC; + if (access & IBV_ACCESS_REMOTE_ATOMIC) + flags |= IONIC_MRF_REMOTE_ATOMIC; - if (access & IBV_ACCESS_MW_BIND) flags |= IONIC_MRF_MW_BIND; + if (access & IBV_ACCESS_MW_BIND) + flags |= IONIC_MRF_MW_BIND; - if (access & IBV_ACCESS_ZERO_BASED) flags |= IONIC_MRF_ZERO_BASED; + if (access & IBV_ACCESS_ZERO_BASED) + flags |= IONIC_MRF_ZERO_BASED; - return flags; + return flags; } /* cqe status indicated in status_length field when err bit is set */ enum ionic_status { - IONIC_STS_OK, - IONIC_STS_LOCAL_LEN_ERR, - IONIC_STS_LOCAL_QP_OPER_ERR, - IONIC_STS_LOCAL_PROT_ERR, - IONIC_STS_WQE_FLUSHED_ERR, - IONIC_STS_MEM_MGMT_OPER_ERR, - IONIC_STS_BAD_RESP_ERR, - IONIC_STS_LOCAL_ACC_ERR, - IONIC_STS_REMOTE_INV_REQ_ERR, - IONIC_STS_REMOTE_ACC_ERR, - IONIC_STS_REMOTE_OPER_ERR, - IONIC_STS_RETRY_EXCEEDED, - IONIC_STS_RNR_RETRY_EXCEEDED, - IONIC_STS_XRC_VIO_ERR, + IONIC_STS_OK, + IONIC_STS_LOCAL_LEN_ERR, + IONIC_STS_LOCAL_QP_OPER_ERR, + IONIC_STS_LOCAL_PROT_ERR, + IONIC_STS_WQE_FLUSHED_ERR, + IONIC_STS_MEM_MGMT_OPER_ERR, + IONIC_STS_BAD_RESP_ERR, + IONIC_STS_LOCAL_ACC_ERR, + IONIC_STS_REMOTE_INV_REQ_ERR, + IONIC_STS_REMOTE_ACC_ERR, + IONIC_STS_REMOTE_OPER_ERR, + IONIC_STS_RETRY_EXCEEDED, + IONIC_STS_RNR_RETRY_EXCEEDED, + IONIC_STS_XRC_VIO_ERR, }; -static inline int ionic_to_ibv_status(int sts) { - switch (sts) { - case IONIC_STS_OK: - return IBV_WC_SUCCESS; - case IONIC_STS_LOCAL_LEN_ERR: - return IBV_WC_LOC_LEN_ERR; - case IONIC_STS_LOCAL_QP_OPER_ERR: - return IBV_WC_LOC_QP_OP_ERR; - case IONIC_STS_LOCAL_PROT_ERR: - return IBV_WC_LOC_PROT_ERR; - case IONIC_STS_WQE_FLUSHED_ERR: - return IBV_WC_WR_FLUSH_ERR; - case IONIC_STS_MEM_MGMT_OPER_ERR: - return IBV_WC_MW_BIND_ERR; - case IONIC_STS_BAD_RESP_ERR: - return IBV_WC_BAD_RESP_ERR; - case IONIC_STS_LOCAL_ACC_ERR: - return IBV_WC_LOC_ACCESS_ERR; - case IONIC_STS_REMOTE_INV_REQ_ERR: - return IBV_WC_REM_INV_REQ_ERR; - case IONIC_STS_REMOTE_ACC_ERR: - return IBV_WC_REM_ACCESS_ERR; - case IONIC_STS_REMOTE_OPER_ERR: - return IBV_WC_REM_OP_ERR; - case IONIC_STS_RETRY_EXCEEDED: - return IBV_WC_RETRY_EXC_ERR; - case IONIC_STS_RNR_RETRY_EXCEEDED: - return IBV_WC_RNR_RETRY_EXC_ERR; - case IONIC_STS_XRC_VIO_ERR: - default: - return IBV_WC_GENERAL_ERR; - } +static inline int ionic_to_ibv_status(int sts) +{ + switch (sts) { + case IONIC_STS_OK: + return IBV_WC_SUCCESS; + case IONIC_STS_LOCAL_LEN_ERR: + return IBV_WC_LOC_LEN_ERR; + case IONIC_STS_LOCAL_QP_OPER_ERR: + return IBV_WC_LOC_QP_OP_ERR; + case IONIC_STS_LOCAL_PROT_ERR: + return IBV_WC_LOC_PROT_ERR; + case IONIC_STS_WQE_FLUSHED_ERR: + return IBV_WC_WR_FLUSH_ERR; + case IONIC_STS_MEM_MGMT_OPER_ERR: + return IBV_WC_MW_BIND_ERR; + case IONIC_STS_BAD_RESP_ERR: + return IBV_WC_BAD_RESP_ERR; + case IONIC_STS_LOCAL_ACC_ERR: + return IBV_WC_LOC_ACCESS_ERR; + case IONIC_STS_REMOTE_INV_REQ_ERR: + return IBV_WC_REM_INV_REQ_ERR; + case IONIC_STS_REMOTE_ACC_ERR: + return IBV_WC_REM_ACCESS_ERR; + case IONIC_STS_REMOTE_OPER_ERR: + return IBV_WC_REM_OP_ERR; + case IONIC_STS_RETRY_EXCEEDED: + return IBV_WC_RETRY_EXC_ERR; + case IONIC_STS_RNR_RETRY_EXCEEDED: + return IBV_WC_RNR_RETRY_EXC_ERR; + case IONIC_STS_XRC_VIO_ERR: + default: + return IBV_WC_GENERAL_ERR; + } } /* fw abi v1 */ /* data payload part of v1 wqe */ union ionic_v1_pld { - struct ionic_sge sgl[2]; - __be32 spec32[8]; - __be16 spec16[16]; - __u8 data[32]; + struct ionic_sge sgl[2]; + __be32 spec32[8]; + __be16 spec16[16]; + __u8 data[32]; +}; + +struct ionic_v1_cqe_send { + __u8 rsvd[4]; + __be32 msg_msn; + __u8 rsvd2[8]; + __le64 npg_wqe_idx_timestamp; +}; + +struct ionic_v1_cqe_recv { + __le64 wqe_idx_timestamp; + __be32 src_qpn_op; + __u8 src_mac[6]; + __be16 vlan_tag; + __be32 imm_data_rkey; +}; + +struct ionic_v1_cqe_rcqe { + __be64 wqe_idx_timestamp; + __u8 rsvd[8]; + __be32 seq_op_flags; + __be32 imm_data_rkey; }; /* completion queue v1 cqe */ struct ionic_v1_cqe { - union { - struct { - __le64 wqe_idx_timestamp; - __be32 src_qpn_op; - __u8 src_mac[6]; - __be16 vlan_tag; - __be32 imm_data_rkey; - } recv; - struct { - __u8 rsvd[4]; - __be32 msg_msn; - __u8 rsvd2[8]; - __le64 npg_wqe_idx_timestamp; - } send; - }; - __be32 status_length; - __be32 qid_type_flags; + union { + struct ionic_v1_cqe_send send; + struct ionic_v1_cqe_recv recv; + struct ionic_v1_cqe_rcqe rcqe; + }; + __be32 status_length; + __be32 qid_type_flags; }; /* bits for cqe wqe_idx and timestamp */ enum ionic_v1_cqe_wqe_idx_timestamp_bits { - IONIC_V1_CQE_WQE_IDX_MASK = 0xffff, - IONIC_V1_CQE_TIMESTAMP_SHIFT = 16, + IONIC_V1_CQE_WQE_IDX_MASK = 0xffff, + IONIC_V1_CQE_TIMESTAMP_SHIFT = 16, }; +/* bits for rcqe seq_op_flags */ +enum ionic_v1_cqe_rcqe_op_flag_bits { + IONIC_V1_CQE_RCQE_SEQ_MASK = 0xffffff, + IONIC_V1_CQE_RCQE_FLAG_V = BIT(24), + IONIC_V1_CQE_RCQE_FLAG_I = BIT(25), + IONIC_V1_CQE_RCQE_OP_SHIFT = 28, +}; + +static inline uint32_t ionic_v1_rcqe_seq(uint32_t seq_opf) +{ + return seq_opf & IONIC_V1_CQE_RCQE_SEQ_MASK; +} + +static inline uint8_t ionic_v1_rcqe_op(uint32_t seq_opf) +{ + return seq_opf >> IONIC_V1_CQE_RCQE_OP_SHIFT; +} + +static inline bool ionic_v1_rcqe_valid(uint32_t seq_opf) +{ + return seq_opf & IONIC_V1_CQE_RCQE_FLAG_V; +} + +static inline bool ionic_v1_rcqe_ready(uint32_t seq_opf) +{ + return seq_opf & IONIC_V1_CQE_RCQE_FLAG_I; +} + /* bits for cqe recv */ enum ionic_v1_cqe_src_qpn_bits { - IONIC_V1_CQE_RECV_QPN_MASK = 0xffffff, - IONIC_V1_CQE_RECV_OP_SHIFT = 24, - - /* MASK could be 0x3, but need 0x1f for makeshift values: - * OP_TYPE_RDMA_OPER_WITH_IMM, OP_TYPE_SEND_RCVD - */ - IONIC_V1_CQE_RECV_OP_MASK = 0x1f, - IONIC_V1_CQE_RECV_OP_SEND = 0, - IONIC_V1_CQE_RECV_OP_SEND_INV = 1, - IONIC_V1_CQE_RECV_OP_SEND_IMM = 2, - IONIC_V1_CQE_RECV_OP_RDMA_IMM = 3, - - IONIC_V1_CQE_RECV_IS_IPV4 = BIT(7 + IONIC_V1_CQE_RECV_OP_SHIFT), - IONIC_V1_CQE_RECV_IS_VLAN = BIT(6 + IONIC_V1_CQE_RECV_OP_SHIFT), + IONIC_V1_CQE_RECV_QPN_MASK = 0xffffff, + IONIC_V1_CQE_RECV_OP_SHIFT = 24, + + /* MASK could be 0x3, but need 0x1f for makeshift values: + * OP_TYPE_RDMA_OPER_WITH_IMM, OP_TYPE_SEND_RCVD + */ + IONIC_V1_CQE_RECV_OP_MASK = 0x1f, + IONIC_V1_CQE_RECV_OP_SEND = 0, + IONIC_V1_CQE_RECV_OP_SEND_INV = 1, + IONIC_V1_CQE_RECV_OP_SEND_IMM = 2, + IONIC_V1_CQE_RECV_OP_RDMA_IMM = 3, + + IONIC_V1_CQE_RECV_IS_IPV4 = BIT(7 + IONIC_V1_CQE_RECV_OP_SHIFT), + IONIC_V1_CQE_RECV_IS_VLAN = BIT(6 + IONIC_V1_CQE_RECV_OP_SHIFT), }; /* bits for cqe qid_type_flags */ enum ionic_v1_cqe_qtf_bits { - IONIC_V1_CQE_COLOR = BIT(0), - IONIC_V1_CQE_ERROR = BIT(1), - IONIC_V1_CQE_TYPE_SHIFT = 5, - IONIC_V1_CQE_TYPE_MASK = 0x7, - IONIC_V1_CQE_QID_SHIFT = 8, - - IONIC_V1_CQE_TYPE_RECV = 1, - IONIC_V1_CQE_TYPE_SEND_MSN = 2, - IONIC_V1_CQE_TYPE_SEND_NPG = 3, - IONIC_V1_CQE_TYPE_RECV_INDIR = 4, + IONIC_V1_CQE_COLOR = BIT(0), + IONIC_V1_CQE_ERROR = BIT(1), + IONIC_V1_CQE_TYPE_SHIFT = 5, + IONIC_V1_CQE_TYPE_MASK = 0x7, + IONIC_V1_CQE_QID_SHIFT = 8, + + IONIC_V1_CQE_TYPE_RECV = 1, + IONIC_V1_CQE_TYPE_SEND_MSN = 2, + IONIC_V1_CQE_TYPE_SEND_NPG = 3, + c = 4, }; #if !defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_HCC__) -static inline bool ionic_v1_cqe_color(struct ionic_v1_cqe* cqe) { - return !!(cqe->qid_type_flags & htobe32(IONIC_V1_CQE_COLOR)); +static inline bool ionic_v1_cqe_color(struct ionic_v1_cqe *cqe) +{ + return !!(cqe->qid_type_flags & htobe32(IONIC_V1_CQE_COLOR)); } -static inline bool ionic_v1_cqe_error(struct ionic_v1_cqe* cqe) { - return !!(cqe->qid_type_flags & htobe32(IONIC_V1_CQE_ERROR)); +static inline bool ionic_v1_cqe_error(struct ionic_v1_cqe *cqe) +{ + return !!(cqe->qid_type_flags & htobe32(IONIC_V1_CQE_ERROR)); } -static inline bool ionic_v1_cqe_recv_is_ipv4(struct ionic_v1_cqe* cqe) { - return !!(cqe->recv.src_qpn_op & htobe32(IONIC_V1_CQE_RECV_IS_IPV4)); +static inline bool ionic_v1_cqe_recv_is_ipv4(struct ionic_v1_cqe *cqe) +{ + return !!(cqe->recv.src_qpn_op & + htobe32(IONIC_V1_CQE_RECV_IS_IPV4)); } -static inline bool ionic_v1_cqe_recv_is_vlan(struct ionic_v1_cqe* cqe) { - return !!(cqe->recv.src_qpn_op & htobe32(IONIC_V1_CQE_RECV_IS_VLAN)); +static inline bool ionic_v1_cqe_recv_is_vlan(struct ionic_v1_cqe *cqe) +{ + return !!(cqe->recv.src_qpn_op & + htobe32(IONIC_V1_CQE_RECV_IS_VLAN)); } -static inline void ionic_v1_cqe_clean(struct ionic_v1_cqe* cqe) { - cqe->qid_type_flags |= htobe32(~0u << IONIC_V1_CQE_QID_SHIFT); +static inline void ionic_v1_cqe_clean(struct ionic_v1_cqe *cqe) +{ + cqe->qid_type_flags |= htobe32(~0u << IONIC_V1_CQE_QID_SHIFT); } -static inline uint32_t ionic_v1_cqe_qtf(struct ionic_v1_cqe* cqe) { - return be32toh(cqe->qid_type_flags); +static inline uint32_t ionic_v1_cqe_qtf(struct ionic_v1_cqe *cqe) +{ + return be32toh(cqe->qid_type_flags); } -#endif // !defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_HCC__) +#endif // !defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_HCC__) -static inline uint8_t ionic_v1_cqe_qtf_type(uint32_t qtf) { - return (qtf >> IONIC_V1_CQE_TYPE_SHIFT) & IONIC_V1_CQE_TYPE_MASK; +static inline uint8_t ionic_v1_cqe_qtf_type(uint32_t qtf) +{ + return (qtf >> IONIC_V1_CQE_TYPE_SHIFT) & IONIC_V1_CQE_TYPE_MASK; } -static inline uint32_t ionic_v1_cqe_qtf_qid(uint32_t qtf) { return qtf >> IONIC_V1_CQE_QID_SHIFT; } +static inline uint32_t ionic_v1_cqe_qtf_qid(uint32_t qtf) +{ + return qtf >> IONIC_V1_CQE_QID_SHIFT; +} /* v1 base wqe header */ struct ionic_v1_base_hdr { - __le64 wqe_idx; - __u8 op; - __u8 num_sge_key; - __be16 flags; - __be32 imm_data_key; + __le64 wqe_idx; + __u8 op; + __u8 num_sge_key; + __be16 flags; + __be32 imm_data_key; }; /* v1 receive wqe body */ struct ionic_v1_recv_bdy { - __u8 rsvd[16]; - union ionic_v1_pld pld; + __u8 rsvd[16]; + union ionic_v1_pld pld; }; /* v1 send/rdma wqe body (common, has sgl) */ struct ionic_v1_common_bdy { - union { - struct { - __be32 ah_id; - __be32 dest_qpn; - __be32 dest_qkey; - } send; - struct { - __be32 remote_va_high; - __be32 remote_va_low; - __be32 remote_rkey; - } rdma; - }; - __be32 length; - union ionic_v1_pld pld; + union { + struct { + __be32 ah_id; + __be32 dest_qpn; + __be32 dest_qkey; + } send; + struct { + __be32 remote_va_high; + __be32 remote_va_low; + __be32 remote_rkey; + } rdma; + }; + __be32 length; + union ionic_v1_pld pld; }; /* v1 atomic wqe body */ struct ionic_v1_atomic_bdy { - __be32 remote_va_high; - __be32 remote_va_low; - __be32 remote_rkey; - __be32 swap_add_high; - __be32 swap_add_low; - __be32 compare_high; - __be32 compare_low; - __u8 rsvd[4]; - struct ionic_sge sge; + __be32 remote_va_high; + __be32 remote_va_low; + __be32 remote_rkey; + __be32 swap_add_high; + __be32 swap_add_low; + __be32 compare_high; + __be32 compare_low; + __u8 rsvd[4]; + struct ionic_sge sge; }; /* v2 atomic wqe body */ struct ionic_v2_atomic_bdy { - __be32 remote_va_high; - __be32 remote_va_low; - __be32 remote_rkey; - __be32 swap_add_high; - __be32 swap_add_low; - __be32 compare_high; - __be32 compare_low; - __be32 lkey; - __be64 local_va; - __u8 rsvd_expdb[8]; + __be32 remote_va_high; + __be32 remote_va_low; + __be32 remote_rkey; + __be32 swap_add_high; + __be32 swap_add_low; + __be32 compare_high; + __be32 compare_low; + __be32 lkey; + __be64 local_va; + __u8 rsvd_expdb[8]; }; /* v1 bind mw wqe body */ struct ionic_v1_bind_mw_bdy { - __be64 va; - __be64 length; - __be32 lkey; - __be16 flags; - __u8 rsvd[26]; + __be64 va; + __be64 length; + __be32 lkey; + __be16 flags; + __u8 rsvd[26]; }; /* v1 send/recv wqe */ struct ionic_v1_wqe { - struct ionic_v1_base_hdr base; - union { - struct ionic_v1_recv_bdy recv; - struct ionic_v1_common_bdy common; - struct ionic_v1_atomic_bdy atomic; - struct ionic_v2_atomic_bdy atomic_v2; - struct ionic_v1_bind_mw_bdy bind_mw; - }; + struct ionic_v1_base_hdr base; + union { + struct ionic_v1_recv_bdy recv; + struct ionic_v1_common_bdy common; + struct ionic_v1_atomic_bdy atomic; + struct ionic_v2_atomic_bdy atomic_v2; + struct ionic_v1_bind_mw_bdy bind_mw; + }; }; /* queue pair v1 send opcodes */ enum ionic_v1_op { - IONIC_V1_OP_SEND, - IONIC_V1_OP_SEND_INV, - IONIC_V1_OP_SEND_IMM, - IONIC_V1_OP_RDMA_READ, - IONIC_V1_OP_RDMA_WRITE, - IONIC_V1_OP_RDMA_WRITE_IMM, - IONIC_V1_OP_ATOMIC_CS, - IONIC_V1_OP_ATOMIC_FA, - IONIC_V1_OP_REG_MR, - IONIC_V1_OP_LOCAL_INV, - IONIC_V1_OP_BIND_MW, - - /* flags */ - IONIC_V1_FLAG_FENCE = BIT(0), - IONIC_V1_FLAG_SOL = BIT(1), - IONIC_V1_FLAG_INL = BIT(2), - IONIC_V1_FLAG_SIG = BIT(3), - IONIC_V1_FLAG_COLOR = BIT(4), - - /* flags last four bits for sgl spec format */ - IONIC_V1_FLAG_SPEC32 = (1u << 12), - IONIC_V1_FLAG_SPEC16 = (2u << 12), - IONIC_V1_SPEC_FIRST_SGE = 2, + IONIC_V1_OP_SEND, + IONIC_V1_OP_SEND_INV, + IONIC_V1_OP_SEND_IMM, + IONIC_V1_OP_RDMA_READ, + IONIC_V1_OP_RDMA_WRITE, + IONIC_V1_OP_RDMA_WRITE_IMM, + IONIC_V1_OP_ATOMIC_CS, + IONIC_V1_OP_ATOMIC_FA, + IONIC_V1_OP_REG_MR, + IONIC_V1_OP_LOCAL_INV, + IONIC_V1_OP_BIND_MW, + + /* flags */ + IONIC_V1_FLAG_FENCE = BIT(0), + IONIC_V1_FLAG_SOL = BIT(1), + IONIC_V1_FLAG_INL = BIT(2), + IONIC_V1_FLAG_SIG = BIT(3), + IONIC_V1_FLAG_COLOR = BIT(4), + + /* flags last four bits for sgl spec format */ + IONIC_V1_FLAG_SPEC32 = (1u << 12), + IONIC_V1_FLAG_SPEC16 = (2u << 12), + IONIC_V1_SPEC_FIRST_SGE = 2, }; /* queue pair v2 send opcodes */ enum ionic_v2_op { - IONIC_V2_OPSL_OUT = 0x20, - IONIC_V2_OPSL_IMM = 0x40, - IONIC_V2_OPSL_INV = 0x80, + IONIC_V2_OPSL_OUT = 0x20, + IONIC_V2_OPSL_IMM = 0x40, + IONIC_V2_OPSL_INV = 0x80, - IONIC_V2_OP_SEND = 0x0 | IONIC_V2_OPSL_OUT, - IONIC_V2_OP_SEND_IMM = IONIC_V2_OP_SEND | IONIC_V2_OPSL_IMM, - IONIC_V2_OP_SEND_INV = IONIC_V2_OP_SEND | IONIC_V2_OPSL_INV, + IONIC_V2_OP_SEND = 0x0 | IONIC_V2_OPSL_OUT, + IONIC_V2_OP_SEND_IMM = IONIC_V2_OP_SEND | IONIC_V2_OPSL_IMM, + IONIC_V2_OP_SEND_INV = IONIC_V2_OP_SEND | IONIC_V2_OPSL_INV, - IONIC_V2_OP_RDMA_WRITE = 0x1 | IONIC_V2_OPSL_OUT, - IONIC_V2_OP_RDMA_WRITE_IMM = IONIC_V2_OP_RDMA_WRITE | IONIC_V2_OPSL_IMM, + IONIC_V2_OP_RDMA_WRITE = 0x1 | IONIC_V2_OPSL_OUT, + IONIC_V2_OP_RDMA_WRITE_IMM = IONIC_V2_OP_RDMA_WRITE | IONIC_V2_OPSL_IMM, - IONIC_V2_OP_RDMA_READ = 0x2, + IONIC_V2_OP_RDMA_READ = 0x2, - IONIC_V2_OP_ATOMIC_CS = 0x4, - IONIC_V2_OP_ATOMIC_FA = 0x5, - IONIC_V2_OP_REG_MR = 0x6, - IONIC_V2_OP_LOCAL_INV = 0x7, - IONIC_V2_OP_BIND_MW = 0x8, + IONIC_V2_OP_ATOMIC_CS = 0x4, + IONIC_V2_OP_ATOMIC_FA = 0x5, + IONIC_V2_OP_REG_MR = 0x6, + IONIC_V2_OP_LOCAL_INV = 0x7, + IONIC_V2_OP_BIND_MW = 0x8, }; #if !defined(__cplusplus) -static inline size_t ionic_v1_send_wqe_min_size(int min_sge, int min_data, int spec, bool expdb) { - size_t sz_wqe, sz_sgl, sz_data; +static inline size_t ionic_v1_send_wqe_min_size(int min_sge, int min_data, + int spec, bool expdb) +{ + size_t sz_wqe, sz_sgl, sz_data; - if (spec > IONIC_V1_SPEC_FIRST_SGE) min_sge += IONIC_V1_SPEC_FIRST_SGE; + if (spec > IONIC_V1_SPEC_FIRST_SGE) + min_sge += IONIC_V1_SPEC_FIRST_SGE; - if (expdb) { - min_sge += 1; - min_data += IONIC_EXP_DBELL_SZ; - } + if (expdb) { + min_sge += 1; + min_data += IONIC_EXP_DBELL_SZ; + } - sz_wqe = sizeof(struct ionic_v1_wqe); - sz_sgl = offsetof(struct ionic_v1_wqe, common.pld.sgl[min_sge]); - sz_data = offsetof(struct ionic_v1_wqe, common.pld.data[min_data]); + sz_wqe = sizeof(struct ionic_v1_wqe); + sz_sgl = offsetof(struct ionic_v1_wqe, common.pld.sgl[min_sge]); + sz_data = offsetof(struct ionic_v1_wqe, common.pld.data[min_data]); - if (sz_sgl > sz_wqe) sz_wqe = sz_sgl; + if (sz_sgl > sz_wqe) + sz_wqe = sz_sgl; - if (sz_data > sz_wqe) sz_wqe = sz_data; + if (sz_data > sz_wqe) + sz_wqe = sz_data; - return roundup_pow_of_two(sz_wqe); + return roundup_pow_of_two(sz_wqe); } -static inline int ionic_v1_send_wqe_max_sge(uint8_t stride_log2, int spec, bool expdb) { - struct ionic_v1_wqe* wqe = (void*)0; - struct ionic_sge* sge = (void*)(1ull << stride_log2); - int num_sge = 0; +static inline int ionic_v1_send_wqe_max_sge(uint8_t stride_log2, int spec, bool expdb) +{ + struct ionic_v1_wqe *wqe = (void *)0; + struct ionic_sge *sge = (void *)(1ull << stride_log2); + int num_sge = 0; - if (expdb) sge -= 1; + if (expdb) + sge -= 1; - if (spec > IONIC_V1_SPEC_FIRST_SGE) num_sge = IONIC_V1_SPEC_FIRST_SGE; + if (spec > IONIC_V1_SPEC_FIRST_SGE) + num_sge = IONIC_V1_SPEC_FIRST_SGE; - num_sge = sge - &wqe->common.pld.sgl[num_sge]; + num_sge = sge - &wqe->common.pld.sgl[num_sge]; - if (spec && num_sge > spec) num_sge = spec; + if (spec && num_sge > spec) + num_sge = spec; - return num_sge; + return num_sge; } -static inline int ionic_v1_send_wqe_max_data(uint8_t stride_log2, bool expdb) { - struct ionic_v1_wqe* wqe = (void*)0; - __u8* data = (void*)(1ull << stride_log2); +static inline int ionic_v1_send_wqe_max_data(uint8_t stride_log2, bool expdb) +{ + struct ionic_v1_wqe *wqe = (void *)0; + __u8 *data = (void *)(1ull << stride_log2); - if (expdb) data -= IONIC_EXP_DBELL_SZ; + if (expdb) + data -= IONIC_EXP_DBELL_SZ; - return data - wqe->common.pld.data; + return data - wqe->common.pld.data; } -static inline size_t ionic_v1_recv_wqe_min_size(int min_sge, int spec, bool expdb) { - size_t sz_wqe, sz_sgl; +static inline size_t ionic_v1_recv_wqe_min_size(int min_sge, int spec, bool expdb) +{ + size_t sz_wqe, sz_sgl; - if (spec > IONIC_V1_SPEC_FIRST_SGE) min_sge += IONIC_V1_SPEC_FIRST_SGE; + if (spec > IONIC_V1_SPEC_FIRST_SGE) + min_sge += IONIC_V1_SPEC_FIRST_SGE; - if (expdb) min_sge += 1; + if (expdb) + min_sge += 1; - sz_wqe = sizeof(struct ionic_v1_wqe); - sz_sgl = offsetof(struct ionic_v1_wqe, recv.pld.sgl[min_sge]); + sz_wqe = sizeof(struct ionic_v1_wqe); + sz_sgl = offsetof(struct ionic_v1_wqe, recv.pld.sgl[min_sge]); - if (sz_sgl > sz_wqe) sz_wqe = sz_sgl; + if (sz_sgl > sz_wqe) + sz_wqe = sz_sgl; - return sz_wqe; + return sz_wqe; } -static inline int ionic_v1_recv_wqe_max_sge(uint8_t stride_log2, int spec, bool expdb) { - struct ionic_v1_wqe* wqe = (void*)0; - struct ionic_sge* sge = (void*)(1ull << stride_log2); - int num_sge = 0; +static inline int ionic_v1_recv_wqe_max_sge(uint8_t stride_log2, int spec, bool expdb) +{ + struct ionic_v1_wqe *wqe = (void *)0; + struct ionic_sge *sge = (void *)(1ull << stride_log2); + int num_sge = 0; - if (expdb) sge -= 1; + if (expdb) + sge -= 1; - if (spec > IONIC_V1_SPEC_FIRST_SGE) num_sge = IONIC_V1_SPEC_FIRST_SGE; + if (spec > IONIC_V1_SPEC_FIRST_SGE) + num_sge = IONIC_V1_SPEC_FIRST_SGE; - num_sge = sge - &wqe->recv.pld.sgl[num_sge]; + num_sge = sge - &wqe->recv.pld.sgl[num_sge]; - if (spec && num_sge > spec) num_sge = spec; + if (spec && num_sge > spec) + num_sge = spec; - return num_sge; + return num_sge; } -static inline int ionic_v1_use_spec_sge(int min_sge, int spec) { - if (!spec || min_sge > spec) return 0; +static inline int ionic_v1_use_spec_sge(int min_sge, int spec) +{ + if (!spec || min_sge > spec) + return 0; - if (min_sge <= IONIC_V1_SPEC_FIRST_SGE) return IONIC_V1_SPEC_FIRST_SGE; + if (min_sge <= IONIC_V1_SPEC_FIRST_SGE) + return IONIC_V1_SPEC_FIRST_SGE; - return spec; + return spec; } #define IONIC_RCQ_SIZE 4096 -#define IONIC_RCQ_DEPTH 128 -#define IONIC_RCQ_DEPTH_LOG2 7 -#define IONIC_RCQ_STRIDE_LOG2 4 struct ionic_rcq_hdr { - uint8_t pad[60]; - uint32_t seq_pad; -}; - -struct ionic_rcqe { - uint32_t status_length; - uint32_t imm_data; - uint32_t seq_flags; - uint32_t rsvd; -}; - -enum ionic_rcqe_flag { - IONIC_RCQE_C = BIT(7), - IONIC_RCQE_I = BIT(6), + __be32 seq; + __be32 ack; }; struct ionic_rcq { - struct ionic_rcq_hdr hdr; - struct ionic_rcqe ring[IONIC_RCQ_DEPTH]; + union { + uint8_t bytes[IONIC_RCQ_SIZE]; + struct ionic_rcq_hdr hdr; + }; }; -static inline uint32_t ionic_rcq_hdr_seq(struct ionic_rcq_hdr* hdr) { - return be32toh(hdr->seq_pad) >> 8; -} - -static inline uint32_t ionic_rcqe_seq(struct ionic_rcqe* rcqe) { - return be32toh(rcqe->seq_flags) >> 8; -} - -static inline bool ionic_rcqe_color(struct ionic_rcqe* rcqe) { - return !!(rcqe->seq_flags & htobe32(IONIC_RCQE_C)); +static inline uint32_t ionic_rcq_seq(struct ionic_rcq *rcq) +{ + return be32toh(rcq->hdr.seq) & IONIC_V1_CQE_RCQE_SEQ_MASK; } -static inline bool ionic_rcqe_imm(struct ionic_rcqe* rcqe) { - return !!(rcqe->seq_flags & htobe32(IONIC_RCQE_I)); +static inline void ionic_rcq_ack(struct ionic_rcq *rcq, uint32_t ack) +{ + rcq->hdr.ack = htobe32(ack); } -#endif // !defined(__cplusplus) +#endif // !defined(__cplusplus) #endif /* IONIC_FW_H */ From 961c63615111c92fd233c028ba68f54e6300909b Mon Sep 17 00:00:00 2001 From: qizzhang Date: Wed, 6 May 2026 16:45:11 +0800 Subject: [PATCH 02/33] compatiable ccqe/non-ccqe --- python/mori/jit/cache.py | 8 +++-- python/mori/jit/core.py | 35 ++++++++++++++++--- .../transport/rdma/providers/ionic/ionic.cpp | 11 ++++++ 3 files changed, 48 insertions(+), 6 deletions(-) diff --git a/python/mori/jit/cache.py b/python/mori/jit/cache.py index c8ab2ebb..ab4221c4 100644 --- a/python/mori/jit/cache.py +++ b/python/mori/jit/cache.py @@ -62,10 +62,11 @@ def get_cache_dir( profiler: bool = False, *, cov: int | None = None, + ccqe: bool = False, ) -> Path: """Return the cache directory for a specific arch + NIC + content combo. - Structure: /_[_profiler][_cov]// + Structure: /_[_ccqe][_profiler][_cov]// Args: profiler: When True, appends '_profiler' to the directory name so that @@ -74,10 +75,13 @@ def get_cache_dir( included in the directory name to separate bitcode compiled with different ABI versions (e.g. cov5 for Triton, cov6 for FlyDSL). None omits the suffix for backward compatibility. + ccqe: When True, appends '_ccqe' so CCQE and non-CCQE kernels are + cached separately (they differ by -DIONIC_CCQE compile flag). """ content_hash = _hash_tree(source_paths) + ccqe_suffix = "_ccqe" if ccqe else "" profiler_suffix = "_profiler" if profiler else "" cov_suffix = f"_cov{cov}" if cov is not None else "" - d = get_cache_root() / f"{arch}_{nic}{profiler_suffix}{cov_suffix}" / content_hash + d = get_cache_root() / f"{arch}_{nic}{ccqe_suffix}{profiler_suffix}{cov_suffix}" / content_hash d.mkdir(parents=True, exist_ok=True) return d diff --git a/python/mori/jit/core.py b/python/mori/jit/core.py index 825d56f3..d552adb7 100644 --- a/python/mori/jit/core.py +++ b/python/mori/jit/core.py @@ -114,6 +114,7 @@ def _hipcc_device_bc( "-D__HIP_PLATFORM_AMD__", "-DHIP_ENABLE_WARP_SYNC_BUILTINS", *_nic_defines(), + *_ccqe_defines(), *_profiler_defines(), ] for d in include_dirs: @@ -162,6 +163,29 @@ def _verify_bitcode(cfg: BuildConfig, bc_path: Path) -> None: ) +def _has_ionic_ccqe() -> bool: + """Check whether the ionic DV header exposes ionic_dv_create_cq_ex (CCQE support).""" + header = Path("/usr/include/infiniband/ionic_dv.h") + try: + return "ionic_dv_create_cq_ex" in header.read_text() + except OSError: + return False + + +_ccqe_enabled: bool | None = None + + +def is_ccqe_enabled() -> bool: + """Return True if CCQE should be enabled (cached after first call).""" + global _ccqe_enabled + if _ccqe_enabled is None: + _ccqe_enabled = detect_nic_type() == "ionic" and _has_ionic_ccqe() + return _ccqe_enabled + + +def _ccqe_defines() -> list[str]: + return ["-DIONIC_CCQE"] if is_ccqe_enabled() else [] + def _nic_defines() -> list[str]: """Return compiler -D flags for the detected NIC type (device-side macros).""" nic = detect_nic_type() @@ -305,6 +329,7 @@ def _hipcc_genco( "-D__HIP_PLATFORM_AMD__", "-DHIP_ENABLE_WARP_SYNC_BUILTINS", *_nic_defines(), + *_ccqe_defines(), *_profiler_defines(), ] for d in include_dirs: @@ -379,6 +404,7 @@ def compile_genco( cfg = detect_build_config() nic = detect_nic_type() profiler = is_profiler_enabled() + ccqe = is_ccqe_enabled() include_dirs = _collect_include_dirs(mori_root) sub_kernels = _PARALLEL_KERNEL_GROUPS.get(kernel_name) @@ -387,7 +413,7 @@ def compile_genco( mori_root / "src" / "ops" / "kernels", mori_root / "include" / "mori", ] - cache_dir = get_cache_dir(cfg.arch, source_paths, nic, profiler=profiler) + cache_dir = get_cache_dir(cfg.arch, source_paths, nic, profiler=profiler, ccqe=ccqe) hsaco_paths = [cache_dir / f"{k}.hsaco" for k in sub_kernels] if all(p.is_file() for p in hsaco_paths): @@ -432,7 +458,7 @@ def compile_genco( raise FileNotFoundError(f"Kernel source not found: {source}") source_paths = [source, mori_root / "include" / "mori"] - cache_dir = get_cache_dir(cfg.arch, source_paths, nic, profiler=profiler) + cache_dir = get_cache_dir(cfg.arch, source_paths, nic, profiler=profiler, ccqe=ccqe) hsaco_path = cache_dir / f"{kernel_name}.hsaco" if hsaco_path.is_file(): @@ -448,7 +474,7 @@ def compile_genco( nic = detect_nic_type() print( f"[mori-jit] Compiling {kernel_name} for {cfg.arch} " - f"(nic={nic}, profiler={profiler}) ..." + f"(nic={nic}, ccqe={ccqe}, profiler={profiler}) ..." ) _hipcc_genco(cfg, source, include_dirs, hsaco_path) print(f"[mori-jit] Cached: {hsaco_path}") @@ -476,12 +502,13 @@ def ensure_bitcode(*, cov: int = 5) -> str: nic = detect_nic_type() profiler = is_profiler_enabled() + ccqe = is_ccqe_enabled() source_paths = [ mori_root / "src" / "shmem" / "shmem_device_api_wrapper.cpp", mori_root / "include" / "mori" / "shmem", mori_root / "include" / "mori" / "core", ] - cache_dir = get_cache_dir(cfg.arch, source_paths, nic, profiler=profiler, cov=cov) + cache_dir = get_cache_dir(cfg.arch, source_paths, nic, profiler=profiler, cov=cov, ccqe=ccqe) bc_path = cache_dir / _BC_FILENAME if bc_path.is_file(): diff --git a/src/application/transport/rdma/providers/ionic/ionic.cpp b/src/application/transport/rdma/providers/ionic/ionic.cpp index e5fea823..96b812b5 100644 --- a/src/application/transport/rdma/providers/ionic/ionic.cpp +++ b/src/application/transport/rdma/providers/ionic/ionic.cpp @@ -67,7 +67,18 @@ IonicCqContainer::IonicCqContainer(ibv_context* context, const RdmaEndpointConfi cq_attr.comp_mask = IBV_CQ_INIT_ATTR_MASK_PD; cq_attr.parent_domain = pd; +#ifdef IONIC_CCQE + struct ionic_cq_init_attr_ex ionic_cq_attr; + memset(&ionic_cq_attr, 0, sizeof(struct ionic_cq_init_attr_ex)); + ionic_cq_attr.comp_mask = IONIC_CQ_INIT_ATTR_MASK_FLAGS; + ionic_cq_attr.flags = IONIC_CQ_INIT_ATTR_CCQE; + cq_attr.cqe = 1; + cq_ex = ionic_dv_create_cq_ex(context, &cq_attr, &ionic_cq_attr); +#else + cq_attr.cqe = cqeNum * 2; // from rocshmem, send&recv? cq_ex = ibv_create_cq_ex(context, &cq_attr); +#endif + assert(cq_ex); cq = ibv_cq_ex_to_cq(cq_ex); assert(cq); From b07315ffbc5e72883f2bf29afa652e90bf756220 Mon Sep 17 00:00:00 2001 From: qizzhang Date: Fri, 8 May 2026 10:37:35 +0800 Subject: [PATCH 03/33] dlopen ionic_dv_create_cq_ex --- .../mori/application/transport/rdma/providers/dv_loader.hpp | 6 +++++- .../mori/core/transport/rdma/providers/ionic/ionic_defs.hpp | 4 ++-- src/application/transport/rdma/providers/ionic/ionic.cpp | 2 +- 3 files changed, 8 insertions(+), 4 deletions(-) diff --git a/include/mori/application/transport/rdma/providers/dv_loader.hpp b/include/mori/application/transport/rdma/providers/dv_loader.hpp index d008effe..b6dfde87 100644 --- a/include/mori/application/transport/rdma/providers/dv_loader.hpp +++ b/include/mori/application/transport/rdma/providers/dv_loader.hpp @@ -181,6 +181,8 @@ struct IonicDvApi { using pd_set_sqcmb_t = int (*)(struct ibv_pd*, bool, bool, bool); using pd_set_rqcmb_t = int (*)(struct ibv_pd*, bool, bool, bool); using pd_set_udma_mask_t = int (*)(struct ibv_pd*, uint32_t); + using create_cq_ex_t = struct ibv_cq_ex* (*)(struct ibv_context*, struct ibv_cq_init_attr_ex*, + struct ionic_cq_init_attr_ex*); get_ctx_t get_ctx = nullptr; qp_get_udma_idx_t qp_get_udma_idx = nullptr; @@ -189,6 +191,7 @@ struct IonicDvApi { pd_set_sqcmb_t pd_set_sqcmb = nullptr; pd_set_rqcmb_t pd_set_rqcmb = nullptr; pd_set_udma_mask_t pd_set_udma_mask = nullptr; + create_cq_ex_t create_cq_ex = nullptr; void* handle = nullptr; @@ -203,9 +206,10 @@ struct IonicDvApi { pd_set_sqcmb = (pd_set_sqcmb_t)DvLoadSymbol(handle, "ionic_dv_pd_set_sqcmb"); pd_set_rqcmb = (pd_set_rqcmb_t)DvLoadSymbol(handle, "ionic_dv_pd_set_rqcmb"); pd_set_udma_mask = (pd_set_udma_mask_t)DvLoadSymbol(handle, "ionic_dv_pd_set_udma_mask"); + create_cq_ex = (create_cq_ex_t)DvLoadSymbol(handle, "ionic_dv_create_cq_ex"); return get_ctx && qp_get_udma_idx && get_cq && get_qp && pd_set_sqcmb && pd_set_rqcmb && - pd_set_udma_mask; + pd_set_udma_mask && create_cq_ex; } static IonicDvApi& Instance() { diff --git a/include/mori/core/transport/rdma/providers/ionic/ionic_defs.hpp b/include/mori/core/transport/rdma/providers/ionic/ionic_defs.hpp index 2e79e875..f6c8562a 100644 --- a/include/mori/core/transport/rdma/providers/ionic/ionic_defs.hpp +++ b/include/mori/core/transport/rdma/providers/ionic/ionic_defs.hpp @@ -26,7 +26,7 @@ namespace core { #define QUEUE_SIZE 1 #define MAX_INLINE_SIZE 32 -// #define IONIC_CCQE 1 -#undef IONIC_CCQE +// // #define IONIC_CCQE 1 +// #undef IONIC_CCQE } // namespace core } // namespace mori diff --git a/src/application/transport/rdma/providers/ionic/ionic.cpp b/src/application/transport/rdma/providers/ionic/ionic.cpp index 96b812b5..1cabe590 100644 --- a/src/application/transport/rdma/providers/ionic/ionic.cpp +++ b/src/application/transport/rdma/providers/ionic/ionic.cpp @@ -73,7 +73,7 @@ IonicCqContainer::IonicCqContainer(ibv_context* context, const RdmaEndpointConfi ionic_cq_attr.comp_mask = IONIC_CQ_INIT_ATTR_MASK_FLAGS; ionic_cq_attr.flags = IONIC_CQ_INIT_ATTR_CCQE; cq_attr.cqe = 1; - cq_ex = ionic_dv_create_cq_ex(context, &cq_attr, &ionic_cq_attr); + cq_ex = IonicDvApi::Instance().create_cq_ex(context, &cq_attr, &ionic_cq_attr); #else cq_attr.cqe = cqeNum * 2; // from rocshmem, send&recv? cq_ex = ibv_create_cq_ex(context, &cq_attr); From 24ac86dcc2b9641c9b471830d8ac6bd1e3c1cf77 Mon Sep 17 00:00:00 2001 From: qizzhang Date: Fri, 8 May 2026 10:56:36 +0800 Subject: [PATCH 04/33] fix ccqe check logic --- python/mori/jit/core.py | 11 ++++++++--- 1 file changed, 8 insertions(+), 3 deletions(-) diff --git a/python/mori/jit/core.py b/python/mori/jit/core.py index d552adb7..2f1f781b 100644 --- a/python/mori/jit/core.py +++ b/python/mori/jit/core.py @@ -164,10 +164,14 @@ def _verify_bitcode(cfg: BuildConfig, bc_path: Path) -> None: def _has_ionic_ccqe() -> bool: - """Check whether the ionic DV header exposes ionic_dv_create_cq_ex (CCQE support).""" - header = Path("/usr/include/infiniband/ionic_dv.h") + """Check whether the ionic driver supports CCQE by probing the runtime library symbol.""" + import ctypes, ctypes.util + lib_name = ctypes.util.find_library("ionic") + if lib_name is None: + return False try: - return "ionic_dv_create_cq_ex" in header.read_text() + lib = ctypes.CDLL(lib_name) + return hasattr(lib, "ionic_dv_create_cq_ex") except OSError: return False @@ -180,6 +184,7 @@ def is_ccqe_enabled() -> bool: global _ccqe_enabled if _ccqe_enabled is None: _ccqe_enabled = detect_nic_type() == "ionic" and _has_ionic_ccqe() + print("xxxxxxxxxxxxxxxxxx CCQE enabled XXXXXXXXXXXXXXXXXXXXX") return _ccqe_enabled From 538f42cca263bb5445423a8a7e502780fe82c0d0 Mon Sep 17 00:00:00 2001 From: qizzhang Date: Fri, 8 May 2026 11:18:36 +0800 Subject: [PATCH 05/33] fix judge --- python/mori/jit/core.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/python/mori/jit/core.py b/python/mori/jit/core.py index 2f1f781b..6fb45226 100644 --- a/python/mori/jit/core.py +++ b/python/mori/jit/core.py @@ -183,7 +183,7 @@ def is_ccqe_enabled() -> bool: """Return True if CCQE should be enabled (cached after first call).""" global _ccqe_enabled if _ccqe_enabled is None: - _ccqe_enabled = detect_nic_type() == "ionic" and _has_ionic_ccqe() + _ccqe_enabled = _has_ionic_ccqe() print("xxxxxxxxxxxxxxxxxx CCQE enabled XXXXXXXXXXXXXXXXXXXXX") return _ccqe_enabled From 76efae7e3856f4048651c64b80a98c2c72178996 Mon Sep 17 00:00:00 2001 From: qizzhang Date: Fri, 8 May 2026 15:33:46 +0800 Subject: [PATCH 06/33] fix print --- python/mori/jit/core.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/python/mori/jit/core.py b/python/mori/jit/core.py index 6fb45226..8ab7b912 100644 --- a/python/mori/jit/core.py +++ b/python/mori/jit/core.py @@ -184,7 +184,7 @@ def is_ccqe_enabled() -> bool: global _ccqe_enabled if _ccqe_enabled is None: _ccqe_enabled = _has_ionic_ccqe() - print("xxxxxxxxxxxxxxxxxx CCQE enabled XXXXXXXXXXXXXXXXXXXXX") + print("xxxxxxxxxxxxxxxxxx CCQE enabled XXXXXXXXXXXXXXXXXXXXX", _ccqe_enabled) return _ccqe_enabled From 838768870963233330584ab4e3ec94beb25a84ef Mon Sep 17 00:00:00 2001 From: qizzhang Date: Fri, 8 May 2026 17:53:50 +0800 Subject: [PATCH 07/33] checkout host ccqe in runtime --- .../transport/rdma/providers/dv_loader.hpp | 3 +- .../transport/rdma/providers/ionic/ionic.cpp | 33 +++++++++---------- 2 files changed, 17 insertions(+), 19 deletions(-) diff --git a/include/mori/application/transport/rdma/providers/dv_loader.hpp b/include/mori/application/transport/rdma/providers/dv_loader.hpp index b6dfde87..fe91b35e 100644 --- a/include/mori/application/transport/rdma/providers/dv_loader.hpp +++ b/include/mori/application/transport/rdma/providers/dv_loader.hpp @@ -208,8 +208,9 @@ struct IonicDvApi { pd_set_udma_mask = (pd_set_udma_mask_t)DvLoadSymbol(handle, "ionic_dv_pd_set_udma_mask"); create_cq_ex = (create_cq_ex_t)DvLoadSymbol(handle, "ionic_dv_create_cq_ex"); + // create_cq_ex is optional: nullptr means CCQE not supported by this driver version return get_ctx && qp_get_udma_idx && get_cq && get_qp && pd_set_sqcmb && pd_set_rqcmb && - pd_set_udma_mask && create_cq_ex; + pd_set_udma_mask; } static IonicDvApi& Instance() { diff --git a/src/application/transport/rdma/providers/ionic/ionic.cpp b/src/application/transport/rdma/providers/ionic/ionic.cpp index 1cabe590..64df880b 100644 --- a/src/application/transport/rdma/providers/ionic/ionic.cpp +++ b/src/application/transport/rdma/providers/ionic/ionic.cpp @@ -52,14 +52,9 @@ IonicCqContainer::IonicCqContainer(ibv_context* context, const RdmaEndpointConfi cqeNum = config.maxCqeNum; + const bool ccqe_enabled = IonicDvApi::Instance().create_cq_ex != nullptr; + memset(&cq_attr, 0, sizeof(struct ibv_cq_init_attr_ex)); -#ifdef IONIC_CCQE - cq_attr.cqe = 0; - MORI_APP_TRACE("cqe mode: ccqe mode"); -#else - cq_attr.cqe = cqeNum * 2; // from rocshmem, send&recv? - MORI_APP_TRACE("cqe mode: normal mode"); -#endif cq_attr.cq_context = nullptr; cq_attr.channel = nullptr; cq_attr.comp_vector = 0; @@ -67,17 +62,19 @@ IonicCqContainer::IonicCqContainer(ibv_context* context, const RdmaEndpointConfi cq_attr.comp_mask = IBV_CQ_INIT_ATTR_MASK_PD; cq_attr.parent_domain = pd; -#ifdef IONIC_CCQE - struct ionic_cq_init_attr_ex ionic_cq_attr; - memset(&ionic_cq_attr, 0, sizeof(struct ionic_cq_init_attr_ex)); - ionic_cq_attr.comp_mask = IONIC_CQ_INIT_ATTR_MASK_FLAGS; - ionic_cq_attr.flags = IONIC_CQ_INIT_ATTR_CCQE; - cq_attr.cqe = 1; - cq_ex = IonicDvApi::Instance().create_cq_ex(context, &cq_attr, &ionic_cq_attr); -#else - cq_attr.cqe = cqeNum * 2; // from rocshmem, send&recv? - cq_ex = ibv_create_cq_ex(context, &cq_attr); -#endif + if (ccqe_enabled) { + MORI_APP_TRACE("cqe mode: ccqe mode"); + struct ionic_cq_init_attr_ex ionic_cq_attr; + memset(&ionic_cq_attr, 0, sizeof(struct ionic_cq_init_attr_ex)); + ionic_cq_attr.comp_mask = IONIC_CQ_INIT_ATTR_MASK_FLAGS; + ionic_cq_attr.flags = IONIC_CQ_INIT_ATTR_CCQE; + cq_attr.cqe = 1; + cq_ex = IonicDvApi::Instance().create_cq_ex(context, &cq_attr, &ionic_cq_attr); + } else { + MORI_APP_TRACE("cqe mode: normal mode"); + cq_attr.cqe = cqeNum * 2; // from rocshmem, send&recv? + cq_ex = ibv_create_cq_ex(context, &cq_attr); + } assert(cq_ex); cq = ibv_cq_ex_to_cq(cq_ex); From e81f47f86759916365c4f88abc4100e9b206a3cf Mon Sep 17 00:00:00 2001 From: qizzhang Date: Fri, 8 May 2026 18:05:11 +0800 Subject: [PATCH 08/33] format --- .../transport/rdma/providers/ionic/ionic_fw.h | 658 ++++++++---------- python/mori/jit/cache.py | 6 +- python/mori/jit/core.py | 13 +- 3 files changed, 320 insertions(+), 357 deletions(-) diff --git a/include/mori/core/transport/rdma/providers/ionic/ionic_fw.h b/include/mori/core/transport/rdma/providers/ionic/ionic_fw.h index 13000476..51679f7a 100644 --- a/include/mori/core/transport/rdma/providers/ionic/ionic_fw.h +++ b/include/mori/core/transport/rdma/providers/ionic/ionic_fw.h @@ -42,531 +42,483 @@ /* wqe scatter gather element */ struct ionic_sge { - __be64 va; - __be32 len; - __be32 lkey; + __be64 va; + __be32 len; + __be32 lkey; }; /* admin queue mr type */ enum ionic_mr_flags { - /* bits that determine mr access */ - IONIC_MRF_LOCAL_WRITE = BIT(0), - IONIC_MRF_REMOTE_WRITE = BIT(1), - IONIC_MRF_REMOTE_READ = BIT(2), - IONIC_MRF_REMOTE_ATOMIC = BIT(3), - IONIC_MRF_MW_BIND = BIT(4), - IONIC_MRF_ZERO_BASED = BIT(5), - IONIC_MRF_ON_DEMAND = BIT(6), - IONIC_MRF_PB = BIT(7), - IONIC_MRF_ACCESS_MASK = BIT(12) - 1, - - /* bits that determine mr type */ - IONIC_MRF_IS_MW = BIT(14), - IONIC_MRF_INV_EN = BIT(15), - - /* base flags combinations for mr types */ - IONIC_MRF_USER_MR = 0, - IONIC_MRF_PHYS_MR = IONIC_MRF_INV_EN, - IONIC_MRF_MW_1 = IONIC_MRF_IS_MW, - IONIC_MRF_MW_2 = IONIC_MRF_IS_MW | IONIC_MRF_INV_EN, + /* bits that determine mr access */ + IONIC_MRF_LOCAL_WRITE = BIT(0), + IONIC_MRF_REMOTE_WRITE = BIT(1), + IONIC_MRF_REMOTE_READ = BIT(2), + IONIC_MRF_REMOTE_ATOMIC = BIT(3), + IONIC_MRF_MW_BIND = BIT(4), + IONIC_MRF_ZERO_BASED = BIT(5), + IONIC_MRF_ON_DEMAND = BIT(6), + IONIC_MRF_PB = BIT(7), + IONIC_MRF_ACCESS_MASK = BIT(12) - 1, + + /* bits that determine mr type */ + IONIC_MRF_IS_MW = BIT(14), + IONIC_MRF_INV_EN = BIT(15), + + /* base flags combinations for mr types */ + IONIC_MRF_USER_MR = 0, + IONIC_MRF_PHYS_MR = IONIC_MRF_INV_EN, + IONIC_MRF_MW_1 = IONIC_MRF_IS_MW, + IONIC_MRF_MW_2 = IONIC_MRF_IS_MW | IONIC_MRF_INV_EN, }; -static inline int to_ionic_mr_flags(int access) -{ - int flags = 0; +static inline int to_ionic_mr_flags(int access) { + int flags = 0; - if (access & IBV_ACCESS_LOCAL_WRITE) - flags |= IONIC_MRF_LOCAL_WRITE; + if (access & IBV_ACCESS_LOCAL_WRITE) flags |= IONIC_MRF_LOCAL_WRITE; - if (access & IBV_ACCESS_REMOTE_READ) - flags |= IONIC_MRF_REMOTE_READ; + if (access & IBV_ACCESS_REMOTE_READ) flags |= IONIC_MRF_REMOTE_READ; - if (access & IBV_ACCESS_REMOTE_WRITE) - flags |= IONIC_MRF_REMOTE_WRITE; + if (access & IBV_ACCESS_REMOTE_WRITE) flags |= IONIC_MRF_REMOTE_WRITE; - if (access & IBV_ACCESS_REMOTE_ATOMIC) - flags |= IONIC_MRF_REMOTE_ATOMIC; + if (access & IBV_ACCESS_REMOTE_ATOMIC) flags |= IONIC_MRF_REMOTE_ATOMIC; - if (access & IBV_ACCESS_MW_BIND) - flags |= IONIC_MRF_MW_BIND; + if (access & IBV_ACCESS_MW_BIND) flags |= IONIC_MRF_MW_BIND; - if (access & IBV_ACCESS_ZERO_BASED) - flags |= IONIC_MRF_ZERO_BASED; + if (access & IBV_ACCESS_ZERO_BASED) flags |= IONIC_MRF_ZERO_BASED; - return flags; + return flags; } /* cqe status indicated in status_length field when err bit is set */ enum ionic_status { - IONIC_STS_OK, - IONIC_STS_LOCAL_LEN_ERR, - IONIC_STS_LOCAL_QP_OPER_ERR, - IONIC_STS_LOCAL_PROT_ERR, - IONIC_STS_WQE_FLUSHED_ERR, - IONIC_STS_MEM_MGMT_OPER_ERR, - IONIC_STS_BAD_RESP_ERR, - IONIC_STS_LOCAL_ACC_ERR, - IONIC_STS_REMOTE_INV_REQ_ERR, - IONIC_STS_REMOTE_ACC_ERR, - IONIC_STS_REMOTE_OPER_ERR, - IONIC_STS_RETRY_EXCEEDED, - IONIC_STS_RNR_RETRY_EXCEEDED, - IONIC_STS_XRC_VIO_ERR, + IONIC_STS_OK, + IONIC_STS_LOCAL_LEN_ERR, + IONIC_STS_LOCAL_QP_OPER_ERR, + IONIC_STS_LOCAL_PROT_ERR, + IONIC_STS_WQE_FLUSHED_ERR, + IONIC_STS_MEM_MGMT_OPER_ERR, + IONIC_STS_BAD_RESP_ERR, + IONIC_STS_LOCAL_ACC_ERR, + IONIC_STS_REMOTE_INV_REQ_ERR, + IONIC_STS_REMOTE_ACC_ERR, + IONIC_STS_REMOTE_OPER_ERR, + IONIC_STS_RETRY_EXCEEDED, + IONIC_STS_RNR_RETRY_EXCEEDED, + IONIC_STS_XRC_VIO_ERR, }; -static inline int ionic_to_ibv_status(int sts) -{ - switch (sts) { - case IONIC_STS_OK: - return IBV_WC_SUCCESS; - case IONIC_STS_LOCAL_LEN_ERR: - return IBV_WC_LOC_LEN_ERR; - case IONIC_STS_LOCAL_QP_OPER_ERR: - return IBV_WC_LOC_QP_OP_ERR; - case IONIC_STS_LOCAL_PROT_ERR: - return IBV_WC_LOC_PROT_ERR; - case IONIC_STS_WQE_FLUSHED_ERR: - return IBV_WC_WR_FLUSH_ERR; - case IONIC_STS_MEM_MGMT_OPER_ERR: - return IBV_WC_MW_BIND_ERR; - case IONIC_STS_BAD_RESP_ERR: - return IBV_WC_BAD_RESP_ERR; - case IONIC_STS_LOCAL_ACC_ERR: - return IBV_WC_LOC_ACCESS_ERR; - case IONIC_STS_REMOTE_INV_REQ_ERR: - return IBV_WC_REM_INV_REQ_ERR; - case IONIC_STS_REMOTE_ACC_ERR: - return IBV_WC_REM_ACCESS_ERR; - case IONIC_STS_REMOTE_OPER_ERR: - return IBV_WC_REM_OP_ERR; - case IONIC_STS_RETRY_EXCEEDED: - return IBV_WC_RETRY_EXC_ERR; - case IONIC_STS_RNR_RETRY_EXCEEDED: - return IBV_WC_RNR_RETRY_EXC_ERR; - case IONIC_STS_XRC_VIO_ERR: - default: - return IBV_WC_GENERAL_ERR; - } +static inline int ionic_to_ibv_status(int sts) { + switch (sts) { + case IONIC_STS_OK: + return IBV_WC_SUCCESS; + case IONIC_STS_LOCAL_LEN_ERR: + return IBV_WC_LOC_LEN_ERR; + case IONIC_STS_LOCAL_QP_OPER_ERR: + return IBV_WC_LOC_QP_OP_ERR; + case IONIC_STS_LOCAL_PROT_ERR: + return IBV_WC_LOC_PROT_ERR; + case IONIC_STS_WQE_FLUSHED_ERR: + return IBV_WC_WR_FLUSH_ERR; + case IONIC_STS_MEM_MGMT_OPER_ERR: + return IBV_WC_MW_BIND_ERR; + case IONIC_STS_BAD_RESP_ERR: + return IBV_WC_BAD_RESP_ERR; + case IONIC_STS_LOCAL_ACC_ERR: + return IBV_WC_LOC_ACCESS_ERR; + case IONIC_STS_REMOTE_INV_REQ_ERR: + return IBV_WC_REM_INV_REQ_ERR; + case IONIC_STS_REMOTE_ACC_ERR: + return IBV_WC_REM_ACCESS_ERR; + case IONIC_STS_REMOTE_OPER_ERR: + return IBV_WC_REM_OP_ERR; + case IONIC_STS_RETRY_EXCEEDED: + return IBV_WC_RETRY_EXC_ERR; + case IONIC_STS_RNR_RETRY_EXCEEDED: + return IBV_WC_RNR_RETRY_EXC_ERR; + case IONIC_STS_XRC_VIO_ERR: + default: + return IBV_WC_GENERAL_ERR; + } } /* fw abi v1 */ /* data payload part of v1 wqe */ union ionic_v1_pld { - struct ionic_sge sgl[2]; - __be32 spec32[8]; - __be16 spec16[16]; - __u8 data[32]; + struct ionic_sge sgl[2]; + __be32 spec32[8]; + __be16 spec16[16]; + __u8 data[32]; }; struct ionic_v1_cqe_send { - __u8 rsvd[4]; - __be32 msg_msn; - __u8 rsvd2[8]; - __le64 npg_wqe_idx_timestamp; + __u8 rsvd[4]; + __be32 msg_msn; + __u8 rsvd2[8]; + __le64 npg_wqe_idx_timestamp; }; struct ionic_v1_cqe_recv { - __le64 wqe_idx_timestamp; - __be32 src_qpn_op; - __u8 src_mac[6]; - __be16 vlan_tag; - __be32 imm_data_rkey; + __le64 wqe_idx_timestamp; + __be32 src_qpn_op; + __u8 src_mac[6]; + __be16 vlan_tag; + __be32 imm_data_rkey; }; struct ionic_v1_cqe_rcqe { - __be64 wqe_idx_timestamp; - __u8 rsvd[8]; - __be32 seq_op_flags; - __be32 imm_data_rkey; + __be64 wqe_idx_timestamp; + __u8 rsvd[8]; + __be32 seq_op_flags; + __be32 imm_data_rkey; }; /* completion queue v1 cqe */ struct ionic_v1_cqe { - union { - struct ionic_v1_cqe_send send; - struct ionic_v1_cqe_recv recv; - struct ionic_v1_cqe_rcqe rcqe; - }; - __be32 status_length; - __be32 qid_type_flags; + union { + struct ionic_v1_cqe_send send; + struct ionic_v1_cqe_recv recv; + struct ionic_v1_cqe_rcqe rcqe; + }; + __be32 status_length; + __be32 qid_type_flags; }; /* bits for cqe wqe_idx and timestamp */ enum ionic_v1_cqe_wqe_idx_timestamp_bits { - IONIC_V1_CQE_WQE_IDX_MASK = 0xffff, - IONIC_V1_CQE_TIMESTAMP_SHIFT = 16, + IONIC_V1_CQE_WQE_IDX_MASK = 0xffff, + IONIC_V1_CQE_TIMESTAMP_SHIFT = 16, }; /* bits for rcqe seq_op_flags */ enum ionic_v1_cqe_rcqe_op_flag_bits { - IONIC_V1_CQE_RCQE_SEQ_MASK = 0xffffff, - IONIC_V1_CQE_RCQE_FLAG_V = BIT(24), - IONIC_V1_CQE_RCQE_FLAG_I = BIT(25), - IONIC_V1_CQE_RCQE_OP_SHIFT = 28, + IONIC_V1_CQE_RCQE_SEQ_MASK = 0xffffff, + IONIC_V1_CQE_RCQE_FLAG_V = BIT(24), + IONIC_V1_CQE_RCQE_FLAG_I = BIT(25), + IONIC_V1_CQE_RCQE_OP_SHIFT = 28, }; -static inline uint32_t ionic_v1_rcqe_seq(uint32_t seq_opf) -{ - return seq_opf & IONIC_V1_CQE_RCQE_SEQ_MASK; +static inline uint32_t ionic_v1_rcqe_seq(uint32_t seq_opf) { + return seq_opf & IONIC_V1_CQE_RCQE_SEQ_MASK; } -static inline uint8_t ionic_v1_rcqe_op(uint32_t seq_opf) -{ - return seq_opf >> IONIC_V1_CQE_RCQE_OP_SHIFT; +static inline uint8_t ionic_v1_rcqe_op(uint32_t seq_opf) { + return seq_opf >> IONIC_V1_CQE_RCQE_OP_SHIFT; } -static inline bool ionic_v1_rcqe_valid(uint32_t seq_opf) -{ - return seq_opf & IONIC_V1_CQE_RCQE_FLAG_V; +static inline bool ionic_v1_rcqe_valid(uint32_t seq_opf) { + return seq_opf & IONIC_V1_CQE_RCQE_FLAG_V; } -static inline bool ionic_v1_rcqe_ready(uint32_t seq_opf) -{ - return seq_opf & IONIC_V1_CQE_RCQE_FLAG_I; +static inline bool ionic_v1_rcqe_ready(uint32_t seq_opf) { + return seq_opf & IONIC_V1_CQE_RCQE_FLAG_I; } /* bits for cqe recv */ enum ionic_v1_cqe_src_qpn_bits { - IONIC_V1_CQE_RECV_QPN_MASK = 0xffffff, - IONIC_V1_CQE_RECV_OP_SHIFT = 24, - - /* MASK could be 0x3, but need 0x1f for makeshift values: - * OP_TYPE_RDMA_OPER_WITH_IMM, OP_TYPE_SEND_RCVD - */ - IONIC_V1_CQE_RECV_OP_MASK = 0x1f, - IONIC_V1_CQE_RECV_OP_SEND = 0, - IONIC_V1_CQE_RECV_OP_SEND_INV = 1, - IONIC_V1_CQE_RECV_OP_SEND_IMM = 2, - IONIC_V1_CQE_RECV_OP_RDMA_IMM = 3, - - IONIC_V1_CQE_RECV_IS_IPV4 = BIT(7 + IONIC_V1_CQE_RECV_OP_SHIFT), - IONIC_V1_CQE_RECV_IS_VLAN = BIT(6 + IONIC_V1_CQE_RECV_OP_SHIFT), + IONIC_V1_CQE_RECV_QPN_MASK = 0xffffff, + IONIC_V1_CQE_RECV_OP_SHIFT = 24, + + /* MASK could be 0x3, but need 0x1f for makeshift values: + * OP_TYPE_RDMA_OPER_WITH_IMM, OP_TYPE_SEND_RCVD + */ + IONIC_V1_CQE_RECV_OP_MASK = 0x1f, + IONIC_V1_CQE_RECV_OP_SEND = 0, + IONIC_V1_CQE_RECV_OP_SEND_INV = 1, + IONIC_V1_CQE_RECV_OP_SEND_IMM = 2, + IONIC_V1_CQE_RECV_OP_RDMA_IMM = 3, + + IONIC_V1_CQE_RECV_IS_IPV4 = BIT(7 + IONIC_V1_CQE_RECV_OP_SHIFT), + IONIC_V1_CQE_RECV_IS_VLAN = BIT(6 + IONIC_V1_CQE_RECV_OP_SHIFT), }; /* bits for cqe qid_type_flags */ enum ionic_v1_cqe_qtf_bits { - IONIC_V1_CQE_COLOR = BIT(0), - IONIC_V1_CQE_ERROR = BIT(1), - IONIC_V1_CQE_TYPE_SHIFT = 5, - IONIC_V1_CQE_TYPE_MASK = 0x7, - IONIC_V1_CQE_QID_SHIFT = 8, - - IONIC_V1_CQE_TYPE_RECV = 1, - IONIC_V1_CQE_TYPE_SEND_MSN = 2, - IONIC_V1_CQE_TYPE_SEND_NPG = 3, - c = 4, + IONIC_V1_CQE_COLOR = BIT(0), + IONIC_V1_CQE_ERROR = BIT(1), + IONIC_V1_CQE_TYPE_SHIFT = 5, + IONIC_V1_CQE_TYPE_MASK = 0x7, + IONIC_V1_CQE_QID_SHIFT = 8, + + IONIC_V1_CQE_TYPE_RECV = 1, + IONIC_V1_CQE_TYPE_SEND_MSN = 2, + IONIC_V1_CQE_TYPE_SEND_NPG = 3, + c = 4, }; #if !defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_HCC__) -static inline bool ionic_v1_cqe_color(struct ionic_v1_cqe *cqe) -{ - return !!(cqe->qid_type_flags & htobe32(IONIC_V1_CQE_COLOR)); +static inline bool ionic_v1_cqe_color(struct ionic_v1_cqe* cqe) { + return !!(cqe->qid_type_flags & htobe32(IONIC_V1_CQE_COLOR)); } -static inline bool ionic_v1_cqe_error(struct ionic_v1_cqe *cqe) -{ - return !!(cqe->qid_type_flags & htobe32(IONIC_V1_CQE_ERROR)); +static inline bool ionic_v1_cqe_error(struct ionic_v1_cqe* cqe) { + return !!(cqe->qid_type_flags & htobe32(IONIC_V1_CQE_ERROR)); } -static inline bool ionic_v1_cqe_recv_is_ipv4(struct ionic_v1_cqe *cqe) -{ - return !!(cqe->recv.src_qpn_op & - htobe32(IONIC_V1_CQE_RECV_IS_IPV4)); +static inline bool ionic_v1_cqe_recv_is_ipv4(struct ionic_v1_cqe* cqe) { + return !!(cqe->recv.src_qpn_op & htobe32(IONIC_V1_CQE_RECV_IS_IPV4)); } -static inline bool ionic_v1_cqe_recv_is_vlan(struct ionic_v1_cqe *cqe) -{ - return !!(cqe->recv.src_qpn_op & - htobe32(IONIC_V1_CQE_RECV_IS_VLAN)); +static inline bool ionic_v1_cqe_recv_is_vlan(struct ionic_v1_cqe* cqe) { + return !!(cqe->recv.src_qpn_op & htobe32(IONIC_V1_CQE_RECV_IS_VLAN)); } -static inline void ionic_v1_cqe_clean(struct ionic_v1_cqe *cqe) -{ - cqe->qid_type_flags |= htobe32(~0u << IONIC_V1_CQE_QID_SHIFT); +static inline void ionic_v1_cqe_clean(struct ionic_v1_cqe* cqe) { + cqe->qid_type_flags |= htobe32(~0u << IONIC_V1_CQE_QID_SHIFT); } -static inline uint32_t ionic_v1_cqe_qtf(struct ionic_v1_cqe *cqe) -{ - return be32toh(cqe->qid_type_flags); +static inline uint32_t ionic_v1_cqe_qtf(struct ionic_v1_cqe* cqe) { + return be32toh(cqe->qid_type_flags); } -#endif // !defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_HCC__) +#endif // !defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_HCC__) -static inline uint8_t ionic_v1_cqe_qtf_type(uint32_t qtf) -{ - return (qtf >> IONIC_V1_CQE_TYPE_SHIFT) & IONIC_V1_CQE_TYPE_MASK; +static inline uint8_t ionic_v1_cqe_qtf_type(uint32_t qtf) { + return (qtf >> IONIC_V1_CQE_TYPE_SHIFT) & IONIC_V1_CQE_TYPE_MASK; } -static inline uint32_t ionic_v1_cqe_qtf_qid(uint32_t qtf) -{ - return qtf >> IONIC_V1_CQE_QID_SHIFT; -} +static inline uint32_t ionic_v1_cqe_qtf_qid(uint32_t qtf) { return qtf >> IONIC_V1_CQE_QID_SHIFT; } /* v1 base wqe header */ struct ionic_v1_base_hdr { - __le64 wqe_idx; - __u8 op; - __u8 num_sge_key; - __be16 flags; - __be32 imm_data_key; + __le64 wqe_idx; + __u8 op; + __u8 num_sge_key; + __be16 flags; + __be32 imm_data_key; }; /* v1 receive wqe body */ struct ionic_v1_recv_bdy { - __u8 rsvd[16]; - union ionic_v1_pld pld; + __u8 rsvd[16]; + union ionic_v1_pld pld; }; /* v1 send/rdma wqe body (common, has sgl) */ struct ionic_v1_common_bdy { - union { - struct { - __be32 ah_id; - __be32 dest_qpn; - __be32 dest_qkey; - } send; - struct { - __be32 remote_va_high; - __be32 remote_va_low; - __be32 remote_rkey; - } rdma; - }; - __be32 length; - union ionic_v1_pld pld; + union { + struct { + __be32 ah_id; + __be32 dest_qpn; + __be32 dest_qkey; + } send; + struct { + __be32 remote_va_high; + __be32 remote_va_low; + __be32 remote_rkey; + } rdma; + }; + __be32 length; + union ionic_v1_pld pld; }; /* v1 atomic wqe body */ struct ionic_v1_atomic_bdy { - __be32 remote_va_high; - __be32 remote_va_low; - __be32 remote_rkey; - __be32 swap_add_high; - __be32 swap_add_low; - __be32 compare_high; - __be32 compare_low; - __u8 rsvd[4]; - struct ionic_sge sge; + __be32 remote_va_high; + __be32 remote_va_low; + __be32 remote_rkey; + __be32 swap_add_high; + __be32 swap_add_low; + __be32 compare_high; + __be32 compare_low; + __u8 rsvd[4]; + struct ionic_sge sge; }; /* v2 atomic wqe body */ struct ionic_v2_atomic_bdy { - __be32 remote_va_high; - __be32 remote_va_low; - __be32 remote_rkey; - __be32 swap_add_high; - __be32 swap_add_low; - __be32 compare_high; - __be32 compare_low; - __be32 lkey; - __be64 local_va; - __u8 rsvd_expdb[8]; + __be32 remote_va_high; + __be32 remote_va_low; + __be32 remote_rkey; + __be32 swap_add_high; + __be32 swap_add_low; + __be32 compare_high; + __be32 compare_low; + __be32 lkey; + __be64 local_va; + __u8 rsvd_expdb[8]; }; /* v1 bind mw wqe body */ struct ionic_v1_bind_mw_bdy { - __be64 va; - __be64 length; - __be32 lkey; - __be16 flags; - __u8 rsvd[26]; + __be64 va; + __be64 length; + __be32 lkey; + __be16 flags; + __u8 rsvd[26]; }; /* v1 send/recv wqe */ struct ionic_v1_wqe { - struct ionic_v1_base_hdr base; - union { - struct ionic_v1_recv_bdy recv; - struct ionic_v1_common_bdy common; - struct ionic_v1_atomic_bdy atomic; - struct ionic_v2_atomic_bdy atomic_v2; - struct ionic_v1_bind_mw_bdy bind_mw; - }; + struct ionic_v1_base_hdr base; + union { + struct ionic_v1_recv_bdy recv; + struct ionic_v1_common_bdy common; + struct ionic_v1_atomic_bdy atomic; + struct ionic_v2_atomic_bdy atomic_v2; + struct ionic_v1_bind_mw_bdy bind_mw; + }; }; /* queue pair v1 send opcodes */ enum ionic_v1_op { - IONIC_V1_OP_SEND, - IONIC_V1_OP_SEND_INV, - IONIC_V1_OP_SEND_IMM, - IONIC_V1_OP_RDMA_READ, - IONIC_V1_OP_RDMA_WRITE, - IONIC_V1_OP_RDMA_WRITE_IMM, - IONIC_V1_OP_ATOMIC_CS, - IONIC_V1_OP_ATOMIC_FA, - IONIC_V1_OP_REG_MR, - IONIC_V1_OP_LOCAL_INV, - IONIC_V1_OP_BIND_MW, - - /* flags */ - IONIC_V1_FLAG_FENCE = BIT(0), - IONIC_V1_FLAG_SOL = BIT(1), - IONIC_V1_FLAG_INL = BIT(2), - IONIC_V1_FLAG_SIG = BIT(3), - IONIC_V1_FLAG_COLOR = BIT(4), - - /* flags last four bits for sgl spec format */ - IONIC_V1_FLAG_SPEC32 = (1u << 12), - IONIC_V1_FLAG_SPEC16 = (2u << 12), - IONIC_V1_SPEC_FIRST_SGE = 2, + IONIC_V1_OP_SEND, + IONIC_V1_OP_SEND_INV, + IONIC_V1_OP_SEND_IMM, + IONIC_V1_OP_RDMA_READ, + IONIC_V1_OP_RDMA_WRITE, + IONIC_V1_OP_RDMA_WRITE_IMM, + IONIC_V1_OP_ATOMIC_CS, + IONIC_V1_OP_ATOMIC_FA, + IONIC_V1_OP_REG_MR, + IONIC_V1_OP_LOCAL_INV, + IONIC_V1_OP_BIND_MW, + + /* flags */ + IONIC_V1_FLAG_FENCE = BIT(0), + IONIC_V1_FLAG_SOL = BIT(1), + IONIC_V1_FLAG_INL = BIT(2), + IONIC_V1_FLAG_SIG = BIT(3), + IONIC_V1_FLAG_COLOR = BIT(4), + + /* flags last four bits for sgl spec format */ + IONIC_V1_FLAG_SPEC32 = (1u << 12), + IONIC_V1_FLAG_SPEC16 = (2u << 12), + IONIC_V1_SPEC_FIRST_SGE = 2, }; /* queue pair v2 send opcodes */ enum ionic_v2_op { - IONIC_V2_OPSL_OUT = 0x20, - IONIC_V2_OPSL_IMM = 0x40, - IONIC_V2_OPSL_INV = 0x80, + IONIC_V2_OPSL_OUT = 0x20, + IONIC_V2_OPSL_IMM = 0x40, + IONIC_V2_OPSL_INV = 0x80, - IONIC_V2_OP_SEND = 0x0 | IONIC_V2_OPSL_OUT, - IONIC_V2_OP_SEND_IMM = IONIC_V2_OP_SEND | IONIC_V2_OPSL_IMM, - IONIC_V2_OP_SEND_INV = IONIC_V2_OP_SEND | IONIC_V2_OPSL_INV, + IONIC_V2_OP_SEND = 0x0 | IONIC_V2_OPSL_OUT, + IONIC_V2_OP_SEND_IMM = IONIC_V2_OP_SEND | IONIC_V2_OPSL_IMM, + IONIC_V2_OP_SEND_INV = IONIC_V2_OP_SEND | IONIC_V2_OPSL_INV, - IONIC_V2_OP_RDMA_WRITE = 0x1 | IONIC_V2_OPSL_OUT, - IONIC_V2_OP_RDMA_WRITE_IMM = IONIC_V2_OP_RDMA_WRITE | IONIC_V2_OPSL_IMM, + IONIC_V2_OP_RDMA_WRITE = 0x1 | IONIC_V2_OPSL_OUT, + IONIC_V2_OP_RDMA_WRITE_IMM = IONIC_V2_OP_RDMA_WRITE | IONIC_V2_OPSL_IMM, - IONIC_V2_OP_RDMA_READ = 0x2, + IONIC_V2_OP_RDMA_READ = 0x2, - IONIC_V2_OP_ATOMIC_CS = 0x4, - IONIC_V2_OP_ATOMIC_FA = 0x5, - IONIC_V2_OP_REG_MR = 0x6, - IONIC_V2_OP_LOCAL_INV = 0x7, - IONIC_V2_OP_BIND_MW = 0x8, + IONIC_V2_OP_ATOMIC_CS = 0x4, + IONIC_V2_OP_ATOMIC_FA = 0x5, + IONIC_V2_OP_REG_MR = 0x6, + IONIC_V2_OP_LOCAL_INV = 0x7, + IONIC_V2_OP_BIND_MW = 0x8, }; #if !defined(__cplusplus) -static inline size_t ionic_v1_send_wqe_min_size(int min_sge, int min_data, - int spec, bool expdb) -{ - size_t sz_wqe, sz_sgl, sz_data; +static inline size_t ionic_v1_send_wqe_min_size(int min_sge, int min_data, int spec, bool expdb) { + size_t sz_wqe, sz_sgl, sz_data; - if (spec > IONIC_V1_SPEC_FIRST_SGE) - min_sge += IONIC_V1_SPEC_FIRST_SGE; + if (spec > IONIC_V1_SPEC_FIRST_SGE) min_sge += IONIC_V1_SPEC_FIRST_SGE; - if (expdb) { - min_sge += 1; - min_data += IONIC_EXP_DBELL_SZ; - } + if (expdb) { + min_sge += 1; + min_data += IONIC_EXP_DBELL_SZ; + } - sz_wqe = sizeof(struct ionic_v1_wqe); - sz_sgl = offsetof(struct ionic_v1_wqe, common.pld.sgl[min_sge]); - sz_data = offsetof(struct ionic_v1_wqe, common.pld.data[min_data]); + sz_wqe = sizeof(struct ionic_v1_wqe); + sz_sgl = offsetof(struct ionic_v1_wqe, common.pld.sgl[min_sge]); + sz_data = offsetof(struct ionic_v1_wqe, common.pld.data[min_data]); - if (sz_sgl > sz_wqe) - sz_wqe = sz_sgl; + if (sz_sgl > sz_wqe) sz_wqe = sz_sgl; - if (sz_data > sz_wqe) - sz_wqe = sz_data; + if (sz_data > sz_wqe) sz_wqe = sz_data; - return roundup_pow_of_two(sz_wqe); + return roundup_pow_of_two(sz_wqe); } -static inline int ionic_v1_send_wqe_max_sge(uint8_t stride_log2, int spec, bool expdb) -{ - struct ionic_v1_wqe *wqe = (void *)0; - struct ionic_sge *sge = (void *)(1ull << stride_log2); - int num_sge = 0; +static inline int ionic_v1_send_wqe_max_sge(uint8_t stride_log2, int spec, bool expdb) { + struct ionic_v1_wqe* wqe = (void*)0; + struct ionic_sge* sge = (void*)(1ull << stride_log2); + int num_sge = 0; - if (expdb) - sge -= 1; + if (expdb) sge -= 1; - if (spec > IONIC_V1_SPEC_FIRST_SGE) - num_sge = IONIC_V1_SPEC_FIRST_SGE; + if (spec > IONIC_V1_SPEC_FIRST_SGE) num_sge = IONIC_V1_SPEC_FIRST_SGE; - num_sge = sge - &wqe->common.pld.sgl[num_sge]; + num_sge = sge - &wqe->common.pld.sgl[num_sge]; - if (spec && num_sge > spec) - num_sge = spec; + if (spec && num_sge > spec) num_sge = spec; - return num_sge; + return num_sge; } -static inline int ionic_v1_send_wqe_max_data(uint8_t stride_log2, bool expdb) -{ - struct ionic_v1_wqe *wqe = (void *)0; - __u8 *data = (void *)(1ull << stride_log2); +static inline int ionic_v1_send_wqe_max_data(uint8_t stride_log2, bool expdb) { + struct ionic_v1_wqe* wqe = (void*)0; + __u8* data = (void*)(1ull << stride_log2); - if (expdb) - data -= IONIC_EXP_DBELL_SZ; + if (expdb) data -= IONIC_EXP_DBELL_SZ; - return data - wqe->common.pld.data; + return data - wqe->common.pld.data; } -static inline size_t ionic_v1_recv_wqe_min_size(int min_sge, int spec, bool expdb) -{ - size_t sz_wqe, sz_sgl; +static inline size_t ionic_v1_recv_wqe_min_size(int min_sge, int spec, bool expdb) { + size_t sz_wqe, sz_sgl; - if (spec > IONIC_V1_SPEC_FIRST_SGE) - min_sge += IONIC_V1_SPEC_FIRST_SGE; + if (spec > IONIC_V1_SPEC_FIRST_SGE) min_sge += IONIC_V1_SPEC_FIRST_SGE; - if (expdb) - min_sge += 1; + if (expdb) min_sge += 1; - sz_wqe = sizeof(struct ionic_v1_wqe); - sz_sgl = offsetof(struct ionic_v1_wqe, recv.pld.sgl[min_sge]); + sz_wqe = sizeof(struct ionic_v1_wqe); + sz_sgl = offsetof(struct ionic_v1_wqe, recv.pld.sgl[min_sge]); - if (sz_sgl > sz_wqe) - sz_wqe = sz_sgl; + if (sz_sgl > sz_wqe) sz_wqe = sz_sgl; - return sz_wqe; + return sz_wqe; } -static inline int ionic_v1_recv_wqe_max_sge(uint8_t stride_log2, int spec, bool expdb) -{ - struct ionic_v1_wqe *wqe = (void *)0; - struct ionic_sge *sge = (void *)(1ull << stride_log2); - int num_sge = 0; +static inline int ionic_v1_recv_wqe_max_sge(uint8_t stride_log2, int spec, bool expdb) { + struct ionic_v1_wqe* wqe = (void*)0; + struct ionic_sge* sge = (void*)(1ull << stride_log2); + int num_sge = 0; - if (expdb) - sge -= 1; + if (expdb) sge -= 1; - if (spec > IONIC_V1_SPEC_FIRST_SGE) - num_sge = IONIC_V1_SPEC_FIRST_SGE; + if (spec > IONIC_V1_SPEC_FIRST_SGE) num_sge = IONIC_V1_SPEC_FIRST_SGE; - num_sge = sge - &wqe->recv.pld.sgl[num_sge]; + num_sge = sge - &wqe->recv.pld.sgl[num_sge]; - if (spec && num_sge > spec) - num_sge = spec; + if (spec && num_sge > spec) num_sge = spec; - return num_sge; + return num_sge; } -static inline int ionic_v1_use_spec_sge(int min_sge, int spec) -{ - if (!spec || min_sge > spec) - return 0; +static inline int ionic_v1_use_spec_sge(int min_sge, int spec) { + if (!spec || min_sge > spec) return 0; - if (min_sge <= IONIC_V1_SPEC_FIRST_SGE) - return IONIC_V1_SPEC_FIRST_SGE; + if (min_sge <= IONIC_V1_SPEC_FIRST_SGE) return IONIC_V1_SPEC_FIRST_SGE; - return spec; + return spec; } #define IONIC_RCQ_SIZE 4096 struct ionic_rcq_hdr { - __be32 seq; - __be32 ack; + __be32 seq; + __be32 ack; }; struct ionic_rcq { - union { - uint8_t bytes[IONIC_RCQ_SIZE]; - struct ionic_rcq_hdr hdr; - }; + union { + uint8_t bytes[IONIC_RCQ_SIZE]; + struct ionic_rcq_hdr hdr; + }; }; -static inline uint32_t ionic_rcq_seq(struct ionic_rcq *rcq) -{ - return be32toh(rcq->hdr.seq) & IONIC_V1_CQE_RCQE_SEQ_MASK; +static inline uint32_t ionic_rcq_seq(struct ionic_rcq* rcq) { + return be32toh(rcq->hdr.seq) & IONIC_V1_CQE_RCQE_SEQ_MASK; } -static inline void ionic_rcq_ack(struct ionic_rcq *rcq, uint32_t ack) -{ - rcq->hdr.ack = htobe32(ack); +static inline void ionic_rcq_ack(struct ionic_rcq* rcq, uint32_t ack) { + rcq->hdr.ack = htobe32(ack); } -#endif // !defined(__cplusplus) +#endif // !defined(__cplusplus) #endif /* IONIC_FW_H */ diff --git a/python/mori/jit/cache.py b/python/mori/jit/cache.py index ab4221c4..92f9f076 100644 --- a/python/mori/jit/cache.py +++ b/python/mori/jit/cache.py @@ -82,6 +82,10 @@ def get_cache_dir( ccqe_suffix = "_ccqe" if ccqe else "" profiler_suffix = "_profiler" if profiler else "" cov_suffix = f"_cov{cov}" if cov is not None else "" - d = get_cache_root() / f"{arch}_{nic}{ccqe_suffix}{profiler_suffix}{cov_suffix}" / content_hash + d = ( + get_cache_root() + / f"{arch}_{nic}{ccqe_suffix}{profiler_suffix}{cov_suffix}" + / content_hash + ) d.mkdir(parents=True, exist_ok=True) return d diff --git a/python/mori/jit/core.py b/python/mori/jit/core.py index 8ab7b912..070e324e 100644 --- a/python/mori/jit/core.py +++ b/python/mori/jit/core.py @@ -165,7 +165,9 @@ def _verify_bitcode(cfg: BuildConfig, bc_path: Path) -> None: def _has_ionic_ccqe() -> bool: """Check whether the ionic driver supports CCQE by probing the runtime library symbol.""" - import ctypes, ctypes.util + import ctypes + import ctypes.util + lib_name = ctypes.util.find_library("ionic") if lib_name is None: return False @@ -191,6 +193,7 @@ def is_ccqe_enabled() -> bool: def _ccqe_defines() -> list[str]: return ["-DIONIC_CCQE"] if is_ccqe_enabled() else [] + def _nic_defines() -> list[str]: """Return compiler -D flags for the detected NIC type (device-side macros).""" nic = detect_nic_type() @@ -418,7 +421,9 @@ def compile_genco( mori_root / "src" / "ops" / "kernels", mori_root / "include" / "mori", ] - cache_dir = get_cache_dir(cfg.arch, source_paths, nic, profiler=profiler, ccqe=ccqe) + cache_dir = get_cache_dir( + cfg.arch, source_paths, nic, profiler=profiler, ccqe=ccqe + ) hsaco_paths = [cache_dir / f"{k}.hsaco" for k in sub_kernels] if all(p.is_file() for p in hsaco_paths): @@ -513,7 +518,9 @@ def ensure_bitcode(*, cov: int = 5) -> str: mori_root / "include" / "mori" / "shmem", mori_root / "include" / "mori" / "core", ] - cache_dir = get_cache_dir(cfg.arch, source_paths, nic, profiler=profiler, cov=cov, ccqe=ccqe) + cache_dir = get_cache_dir( + cfg.arch, source_paths, nic, profiler=profiler, cov=cov, ccqe=ccqe + ) bc_path = cache_dir / _BC_FILENAME if bc_path.is_file(): From ee8b8f27f39034b1b399d616b099122ab188a4a7 Mon Sep 17 00:00:00 2001 From: qizzhang Date: Fri, 8 May 2026 18:09:23 +0800 Subject: [PATCH 09/33] remove defines --- .../mori/core/transport/rdma/providers/ionic/ionic_defs.hpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/include/mori/core/transport/rdma/providers/ionic/ionic_defs.hpp b/include/mori/core/transport/rdma/providers/ionic/ionic_defs.hpp index f6c8562a..9d3cd9be 100644 --- a/include/mori/core/transport/rdma/providers/ionic/ionic_defs.hpp +++ b/include/mori/core/transport/rdma/providers/ionic/ionic_defs.hpp @@ -26,7 +26,6 @@ namespace core { #define QUEUE_SIZE 1 #define MAX_INLINE_SIZE 32 -// // #define IONIC_CCQE 1 -// #undef IONIC_CCQE + } // namespace core } // namespace mori From a3ec281dd944b6936f1ab5ec0414ebcfc12f40ba Mon Sep 17 00:00:00 2001 From: qizzhang Date: Sat, 9 May 2026 10:48:53 +0800 Subject: [PATCH 10/33] retrigger CI --- include/mori/core/transport/rdma/providers/ionic/ionic_defs.hpp | 1 + 1 file changed, 1 insertion(+) diff --git a/include/mori/core/transport/rdma/providers/ionic/ionic_defs.hpp b/include/mori/core/transport/rdma/providers/ionic/ionic_defs.hpp index 9d3cd9be..3ec5b182 100644 --- a/include/mori/core/transport/rdma/providers/ionic/ionic_defs.hpp +++ b/include/mori/core/transport/rdma/providers/ionic/ionic_defs.hpp @@ -24,6 +24,7 @@ namespace mori { namespace core { +// Remove CCQE Defines #define QUEUE_SIZE 1 #define MAX_INLINE_SIZE 32 From b05ee30c478aa72ce90ebd6e3df62cec30834dbe Mon Sep 17 00:00:00 2001 From: qizzhang Date: Sat, 9 May 2026 16:58:44 +0800 Subject: [PATCH 11/33] restore ccqe defines --- include/mori/core/transport/rdma/providers/ionic/ionic_defs.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/mori/core/transport/rdma/providers/ionic/ionic_defs.hpp b/include/mori/core/transport/rdma/providers/ionic/ionic_defs.hpp index 3ec5b182..a991d3c0 100644 --- a/include/mori/core/transport/rdma/providers/ionic/ionic_defs.hpp +++ b/include/mori/core/transport/rdma/providers/ionic/ionic_defs.hpp @@ -27,6 +27,6 @@ namespace core { // Remove CCQE Defines #define QUEUE_SIZE 1 #define MAX_INLINE_SIZE 32 - +#define IONIC_CCQE 1 } // namespace core } // namespace mori From 96dde32415caa9840c1feb78a97ac1a65e9679bc Mon Sep 17 00:00:00 2001 From: qizzhang Date: Sat, 9 May 2026 17:14:53 +0800 Subject: [PATCH 12/33] debug --- .../mori/core/transport/rdma/providers/ionic/ionic_defs.hpp | 2 +- .../transport/rdma/providers/ionic/ionic_device_primitives.hpp | 3 +++ 2 files changed, 4 insertions(+), 1 deletion(-) diff --git a/include/mori/core/transport/rdma/providers/ionic/ionic_defs.hpp b/include/mori/core/transport/rdma/providers/ionic/ionic_defs.hpp index a991d3c0..08b22fa5 100644 --- a/include/mori/core/transport/rdma/providers/ionic/ionic_defs.hpp +++ b/include/mori/core/transport/rdma/providers/ionic/ionic_defs.hpp @@ -27,6 +27,6 @@ namespace core { // Remove CCQE Defines #define QUEUE_SIZE 1 #define MAX_INLINE_SIZE 32 -#define IONIC_CCQE 1 +// #define IONIC_CCQE 1 } // namespace core } // namespace mori diff --git a/include/mori/core/transport/rdma/providers/ionic/ionic_device_primitives.hpp b/include/mori/core/transport/rdma/providers/ionic/ionic_device_primitives.hpp index f00f0ba9..7af10821 100644 --- a/include/mori/core/transport/rdma/providers/ionic/ionic_device_primitives.hpp +++ b/include/mori/core/transport/rdma/providers/ionic/ionic_device_primitives.hpp @@ -749,6 +749,7 @@ inline __device__ int PollCq(WorkQueueHandle& wqHandle, CompletionQueueHandle& cqHandle, void* cqAddr, uint32_t cqeNum, uint32_t* consIdx, uint16_t* wqeCounter) { + printf("poll cccqe\n"); PollCqOnce2(wqHandle, cqHandle, 1, cqAddr, cqeNum, *consIdx); *wqeCounter = *consIdx; return 0; @@ -759,6 +760,8 @@ inline __device__ int PollCq(WorkQueueHandle& wqHandle, CompletionQueueHandle& cqHandle, void* cqAddr, uint32_t cqeNum, uint32_t* consIdx, uint16_t* wqeCounter) { + printf("poll non cccqe\n"); + uint32_t greed = 10; const uint32_t curConsIdx = *consIdx; uint64_t activemask = GetActiveLaneMask(); From 93ad1d78b90f0af51a5cd34b83d3952b7936544e Mon Sep 17 00:00:00 2001 From: qizzhang Date: Sat, 9 May 2026 17:24:36 +0800 Subject: [PATCH 13/33] dbg2 --- .../rdma/providers/ionic/ionic_device_primitives.hpp | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/include/mori/core/transport/rdma/providers/ionic/ionic_device_primitives.hpp b/include/mori/core/transport/rdma/providers/ionic/ionic_device_primitives.hpp index 7af10821..bd6f62bf 100644 --- a/include/mori/core/transport/rdma/providers/ionic/ionic_device_primitives.hpp +++ b/include/mori/core/transport/rdma/providers/ionic/ionic_device_primitives.hpp @@ -622,6 +622,8 @@ inline __device__ int PollCqOnce2(WorkQueueHandle& wqHandle, CompletionQueueHand volatile struct ionic_v1_cqe* cqe = reinterpret_cast(cqeAddr); uint32_t old, msn = HTOBE32(cqe->send.msg_msn); + printf("poll cccqe2\n"); + consIdx = wqHandle.dbTouchIdx; // MORI_PRINTF("ABH %s:%d here cons %#x msn %#x\n", __func__, __LINE__, consIdx, msn); @@ -639,10 +641,12 @@ inline __device__ int PollCqOnce2(WorkQueueHandle& wqHandle, CompletionQueueHand #else inline __device__ int PollCqOnce2(WorkQueueHandle& wqHandle, CompletionQueueHandle& cqHandle, uint64_t activemask, void* cqeAddr, uint32_t cqeNum, - uint32_t consIdx) { + cduint32_t consIdx) { uint32_t my_logical_lane_id = get_active_lane_num(activemask); uint32_t my_cq_pos = cqHandle.cq_consumer + my_logical_lane_id; + printf("poll nonccqe 2\n"); + uint32_t cqeIdx = my_cq_pos & (cqeNum - 1); char* Addr = reinterpret_cast(cqeAddr) + (cqeIdx * sizeof(struct ionic_v1_cqe)); From b6bd5be0b7eb506c7f2f4faf8940af3c2f2cc8ba Mon Sep 17 00:00:00 2001 From: qizzhang Date: Sat, 9 May 2026 17:28:18 +0800 Subject: [PATCH 14/33] fix --- .../transport/rdma/providers/ionic/ionic_device_primitives.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/mori/core/transport/rdma/providers/ionic/ionic_device_primitives.hpp b/include/mori/core/transport/rdma/providers/ionic/ionic_device_primitives.hpp index bd6f62bf..5afe4172 100644 --- a/include/mori/core/transport/rdma/providers/ionic/ionic_device_primitives.hpp +++ b/include/mori/core/transport/rdma/providers/ionic/ionic_device_primitives.hpp @@ -641,7 +641,7 @@ inline __device__ int PollCqOnce2(WorkQueueHandle& wqHandle, CompletionQueueHand #else inline __device__ int PollCqOnce2(WorkQueueHandle& wqHandle, CompletionQueueHandle& cqHandle, uint64_t activemask, void* cqeAddr, uint32_t cqeNum, - cduint32_t consIdx) { + uint32_t consIdx) { uint32_t my_logical_lane_id = get_active_lane_num(activemask); uint32_t my_cq_pos = cqHandle.cq_consumer + my_logical_lane_id; From c37428cdee0da74ea0d2f58920c893e34a5d992a Mon Sep 17 00:00:00 2001 From: qizzhang Date: Sat, 9 May 2026 17:38:34 +0800 Subject: [PATCH 15/33] remove printf --- .../rdma/providers/ionic/ionic_device_primitives.hpp | 7 ------- 1 file changed, 7 deletions(-) diff --git a/include/mori/core/transport/rdma/providers/ionic/ionic_device_primitives.hpp b/include/mori/core/transport/rdma/providers/ionic/ionic_device_primitives.hpp index 5afe4172..f00f0ba9 100644 --- a/include/mori/core/transport/rdma/providers/ionic/ionic_device_primitives.hpp +++ b/include/mori/core/transport/rdma/providers/ionic/ionic_device_primitives.hpp @@ -622,8 +622,6 @@ inline __device__ int PollCqOnce2(WorkQueueHandle& wqHandle, CompletionQueueHand volatile struct ionic_v1_cqe* cqe = reinterpret_cast(cqeAddr); uint32_t old, msn = HTOBE32(cqe->send.msg_msn); - printf("poll cccqe2\n"); - consIdx = wqHandle.dbTouchIdx; // MORI_PRINTF("ABH %s:%d here cons %#x msn %#x\n", __func__, __LINE__, consIdx, msn); @@ -645,8 +643,6 @@ inline __device__ int PollCqOnce2(WorkQueueHandle& wqHandle, CompletionQueueHand uint32_t my_logical_lane_id = get_active_lane_num(activemask); uint32_t my_cq_pos = cqHandle.cq_consumer + my_logical_lane_id; - printf("poll nonccqe 2\n"); - uint32_t cqeIdx = my_cq_pos & (cqeNum - 1); char* Addr = reinterpret_cast(cqeAddr) + (cqeIdx * sizeof(struct ionic_v1_cqe)); @@ -753,7 +749,6 @@ inline __device__ int PollCq(WorkQueueHandle& wqHandle, CompletionQueueHandle& cqHandle, void* cqAddr, uint32_t cqeNum, uint32_t* consIdx, uint16_t* wqeCounter) { - printf("poll cccqe\n"); PollCqOnce2(wqHandle, cqHandle, 1, cqAddr, cqeNum, *consIdx); *wqeCounter = *consIdx; return 0; @@ -764,8 +759,6 @@ inline __device__ int PollCq(WorkQueueHandle& wqHandle, CompletionQueueHandle& cqHandle, void* cqAddr, uint32_t cqeNum, uint32_t* consIdx, uint16_t* wqeCounter) { - printf("poll non cccqe\n"); - uint32_t greed = 10; const uint32_t curConsIdx = *consIdx; uint64_t activemask = GetActiveLaneMask(); From ae4672db9dc21d66de67f102ad07244900a962a1 Mon Sep 17 00:00:00 2001 From: qizzhang Date: Sat, 9 May 2026 17:53:57 +0800 Subject: [PATCH 16/33] remove print --- python/mori/jit/core.py | 1 - 1 file changed, 1 deletion(-) diff --git a/python/mori/jit/core.py b/python/mori/jit/core.py index 070e324e..38d9d9c5 100644 --- a/python/mori/jit/core.py +++ b/python/mori/jit/core.py @@ -186,7 +186,6 @@ def is_ccqe_enabled() -> bool: global _ccqe_enabled if _ccqe_enabled is None: _ccqe_enabled = _has_ionic_ccqe() - print("xxxxxxxxxxxxxxxxxx CCQE enabled XXXXXXXXXXXXXXXXXXXXX", _ccqe_enabled) return _ccqe_enabled From 008408557c1f58c2615d405fb4dbc77a25575cf6 Mon Sep 17 00:00:00 2001 From: qizzhang Date: Sat, 9 May 2026 17:56:39 +0800 Subject: [PATCH 17/33] dump cmd --- python/mori/jit/core.py | 1 + 1 file changed, 1 insertion(+) diff --git a/python/mori/jit/core.py b/python/mori/jit/core.py index 38d9d9c5..b6d737f2 100644 --- a/python/mori/jit/core.py +++ b/python/mori/jit/core.py @@ -339,6 +339,7 @@ def _hipcc_genco( *_ccqe_defines(), *_profiler_defines(), ] + print("genco cmd", cmd) for d in include_dirs: cmd.extend(["-I", str(d)]) cmd.extend([str(source), "-o", str(output)]) From b6ebc4f891a451a24cfbe271476ac81eb0100c98 Mon Sep 17 00:00:00 2001 From: qizzhang Date: Mon, 11 May 2026 10:59:09 +0800 Subject: [PATCH 18/33] update deect ccqe logic --- python/mori/jit/core.py | 66 +++++++++++++++++++++++++++++++++++++++-- 1 file changed, 64 insertions(+), 2 deletions(-) diff --git a/python/mori/jit/core.py b/python/mori/jit/core.py index b6d737f2..29d40f8b 100644 --- a/python/mori/jit/core.py +++ b/python/mori/jit/core.py @@ -163,7 +163,7 @@ def _verify_bitcode(cfg: BuildConfig, bc_path: Path) -> None: ) -def _has_ionic_ccqe() -> bool: +def _lib_has_ionic_ccqe() -> bool: """Check whether the ionic driver supports CCQE by probing the runtime library symbol.""" import ctypes import ctypes.util @@ -178,6 +178,68 @@ def _has_ionic_ccqe() -> bool: return False +def _parse_ionic_fw_minor(fw_ver: str) -> int | None: + """Parse the build number from an ionic firmware string like '1.117.5-a58'. + + Extracts the numeric part from the suffix after '-', e.g. 'a58' → 58, 'a119' → 119. + Returns None if the string cannot be parsed. + """ + m = re.search(r"-[a-zA-Z]+(\d+)$", fw_ver.strip()) + if m: + return int(m.group(1)) + return None + + +_CCQE_MIN_FW_MINOR = 58 + + +def _is_firmware_support_ccqe(fw_ver: str) -> bool: + """Return True if the firmware version string reports build number >= 58.""" + minor = _parse_ionic_fw_minor(fw_ver) + return minor is not None and minor >= _CCQE_MIN_FW_MINOR + + +def _get_ionic_fw_versions() -> list[str]: + """Return fw_ver strings for every ionic IB device found in sysfs.""" + ib_dir = "/sys/class/infiniband" + versions: list[str] = [] + try: + for dev in os.listdir(ib_dir): + dev_path = os.path.join(ib_dir, dev) + driver_link = os.path.join(dev_path, "device", "driver") + try: + driver_name = os.path.basename(os.readlink(driver_link)) + except OSError: + continue + if driver_name not in ("ionic_rdma", "ionic"): + continue + fw_path = os.path.join(dev_path, "fw_ver") + try: + fw_ver = Path(fw_path).read_text().strip() + versions.append(fw_ver) + except OSError: + pass + except OSError: + pass + return versions + + +def _is_all_ionic_support_ccqe() -> bool: + """Return True only when every ionic device has the same fw version and that version >= 58.""" + versions = _get_ionic_fw_versions() + if not versions: + return False + if len(set(versions)) != 1: + return False + + for ver in versions: + if not _is_firmware_support_ccqe(ver): + print(ver) + return False + + return True + + _ccqe_enabled: bool | None = None @@ -185,7 +247,7 @@ def is_ccqe_enabled() -> bool: """Return True if CCQE should be enabled (cached after first call).""" global _ccqe_enabled if _ccqe_enabled is None: - _ccqe_enabled = _has_ionic_ccqe() + _ccqe_enabled = _lib_has_ionic_ccqe() and _is_all_ionic_support_ccqe() return _ccqe_enabled From 49199bed6a0c091b7397f23700894bf7fa43ec34 Mon Sep 17 00:00:00 2001 From: qizzhang Date: Mon, 11 May 2026 11:06:31 +0800 Subject: [PATCH 19/33] revise print --- python/mori/jit/core.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/python/mori/jit/core.py b/python/mori/jit/core.py index 29d40f8b..f536c809 100644 --- a/python/mori/jit/core.py +++ b/python/mori/jit/core.py @@ -231,10 +231,9 @@ def _is_all_ionic_support_ccqe() -> bool: return False if len(set(versions)) != 1: return False - + for ver in versions: if not _is_firmware_support_ccqe(ver): - print(ver) return False return True @@ -248,6 +247,7 @@ def is_ccqe_enabled() -> bool: global _ccqe_enabled if _ccqe_enabled is None: _ccqe_enabled = _lib_has_ionic_ccqe() and _is_all_ionic_support_ccqe() + print(f"ccqe_enabled {_ccqe_enabled}") return _ccqe_enabled From 1ee356beade503f8d4875fb79db5fcba8988fe9a Mon Sep 17 00:00:00 2001 From: qizzhang Date: Mon, 11 May 2026 11:22:44 +0800 Subject: [PATCH 20/33] revise host logic --- .../transport/rdma/providers/ionic/ionic.cpp | 41 ++++++++++++++++++- 1 file changed, 40 insertions(+), 1 deletion(-) diff --git a/src/application/transport/rdma/providers/ionic/ionic.cpp b/src/application/transport/rdma/providers/ionic/ionic.cpp index 64df880b..ed725145 100644 --- a/src/application/transport/rdma/providers/ionic/ionic.cpp +++ b/src/application/transport/rdma/providers/ionic/ionic.cpp @@ -25,6 +25,8 @@ #include #include +#include +#include #include #include @@ -40,6 +42,43 @@ namespace application { /* Device Attributes */ /* ---------------------------------------------------------------------------------------------- */ +namespace { + +// Minimum firmware build number required for CCQE support (e.g. "1.117.5-a58" → 58). +constexpr int kCcqeMinFwBuild = 58; + +// Read /sys/class/infiniband//fw_ver and return the numeric build suffix +// (the digits after the last '-[letters]' component). Returns -1 on any failure. +int ReadIonicFwBuild(const char* dev_name) { + char path[256]; + snprintf(path, sizeof(path), "/sys/class/infiniband/%s/fw_ver", dev_name); + + FILE* f = fopen(path, "r"); + if (!f) return -1; + + char buf[64] = {}; + fgets(buf, sizeof(buf), f); + fclose(f); + + // Find last '-' then skip letters to reach the build digits. + char* dash = strrchr(buf, '-'); + if (!dash) return -1; + char* p = dash + 1; + while (*p && !isdigit(static_cast(*p))) ++p; + if (!*p) return -1; + return atoi(p); +} + +bool IsCcqeSupported(ibv_context* context) { + if (IonicDvApi::Instance().create_cq_ex == nullptr) return false; + int build = ReadIonicFwBuild(context->device->name); + + MORI_APP_TRACE("dev: %s fw_build %d", context->device->name, build); + return build >= kCcqeMinFwBuild; +} + +} // namespace + /* ---------------------------------------------------------------------------------------------- */ /* IonicCqContainer */ /* ---------------------------------------------------------------------------------------------- */ @@ -52,7 +91,7 @@ IonicCqContainer::IonicCqContainer(ibv_context* context, const RdmaEndpointConfi cqeNum = config.maxCqeNum; - const bool ccqe_enabled = IonicDvApi::Instance().create_cq_ex != nullptr; + const bool ccqe_enabled = IsCcqeSupported(context); memset(&cq_attr, 0, sizeof(struct ibv_cq_init_attr_ex)); cq_attr.cq_context = nullptr; From 5dfbd7cb81862e8d349b2bb86dbfb10d121aa63a Mon Sep 17 00:00:00 2001 From: qizzhang Date: Mon, 11 May 2026 11:28:17 +0800 Subject: [PATCH 21/33] format --- python/mori/jit/core.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/python/mori/jit/core.py b/python/mori/jit/core.py index f536c809..a89a2eb8 100644 --- a/python/mori/jit/core.py +++ b/python/mori/jit/core.py @@ -235,7 +235,7 @@ def _is_all_ionic_support_ccqe() -> bool: for ver in versions: if not _is_firmware_support_ccqe(ver): return False - + return True From 87f5e2ad73c19d849bf4fafa565a3586b54f40f2 Mon Sep 17 00:00:00 2001 From: qizzhang Date: Mon, 11 May 2026 11:33:15 +0800 Subject: [PATCH 22/33] modify log --- python/mori/jit/core.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/python/mori/jit/core.py b/python/mori/jit/core.py index a89a2eb8..96a783f0 100644 --- a/python/mori/jit/core.py +++ b/python/mori/jit/core.py @@ -247,7 +247,7 @@ def is_ccqe_enabled() -> bool: global _ccqe_enabled if _ccqe_enabled is None: _ccqe_enabled = _lib_has_ionic_ccqe() and _is_all_ionic_support_ccqe() - print(f"ccqe_enabled {_ccqe_enabled}") + print(f"Ionic ccqe: {_ccqe_enabled}") return _ccqe_enabled From 5b6a686b2f874c2f3e67b8bfb47a8b1b0471630a Mon Sep 17 00:00:00 2001 From: qizzhang Date: Mon, 11 May 2026 11:36:42 +0800 Subject: [PATCH 23/33] host log --- src/application/transport/rdma/providers/ionic/ionic.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/application/transport/rdma/providers/ionic/ionic.cpp b/src/application/transport/rdma/providers/ionic/ionic.cpp index ed725145..f803bb69 100644 --- a/src/application/transport/rdma/providers/ionic/ionic.cpp +++ b/src/application/transport/rdma/providers/ionic/ionic.cpp @@ -73,7 +73,7 @@ bool IsCcqeSupported(ibv_context* context) { if (IonicDvApi::Instance().create_cq_ex == nullptr) return false; int build = ReadIonicFwBuild(context->device->name); - MORI_APP_TRACE("dev: %s fw_build %d", context->device->name, build); + MORI_APP_TRACE("dev: {} fw_build {}", context->device->name, build); return build >= kCcqeMinFwBuild; } From 127a13941a3250e9cd3d7026e8e7c33b5a8edb75 Mon Sep 17 00:00:00 2001 From: qizzhang Date: Mon, 11 May 2026 11:55:32 +0800 Subject: [PATCH 24/33] ci --- python/mori/jit/core.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/python/mori/jit/core.py b/python/mori/jit/core.py index 96a783f0..9474d3e8 100644 --- a/python/mori/jit/core.py +++ b/python/mori/jit/core.py @@ -247,7 +247,7 @@ def is_ccqe_enabled() -> bool: global _ccqe_enabled if _ccqe_enabled is None: _ccqe_enabled = _lib_has_ionic_ccqe() and _is_all_ionic_support_ccqe() - print(f"Ionic ccqe: {_ccqe_enabled}") + print(f"Ionic _ccqe_enabled: {_ccqe_enabled}") return _ccqe_enabled From 93d600dc99adffa5c86fc7cb632e36188d5e14fb Mon Sep 17 00:00:00 2001 From: qizzhang Date: Mon, 11 May 2026 12:08:07 +0800 Subject: [PATCH 25/33] test ci --- python/mori/jit/core.py | 1 + 1 file changed, 1 insertion(+) diff --git a/python/mori/jit/core.py b/python/mori/jit/core.py index 9474d3e8..1501e336 100644 --- a/python/mori/jit/core.py +++ b/python/mori/jit/core.py @@ -248,6 +248,7 @@ def is_ccqe_enabled() -> bool: if _ccqe_enabled is None: _ccqe_enabled = _lib_has_ionic_ccqe() and _is_all_ionic_support_ccqe() print(f"Ionic _ccqe_enabled: {_ccqe_enabled}") + return _ccqe_enabled From 6c456b0578bd5663f8177eeff0d6d5f41fac6aea Mon Sep 17 00:00:00 2001 From: Qizhou Zhang Date: Mon, 18 May 2026 06:55:46 +0000 Subject: [PATCH 26/33] fix parse of firmware minor version --- python/mori/jit/core.py | 15 ++++++++------- 1 file changed, 8 insertions(+), 7 deletions(-) diff --git a/python/mori/jit/core.py b/python/mori/jit/core.py index 1501e336..8062b1bd 100644 --- a/python/mori/jit/core.py +++ b/python/mori/jit/core.py @@ -184,12 +184,11 @@ def _parse_ionic_fw_minor(fw_ver: str) -> int | None: Extracts the numeric part from the suffix after '-', e.g. 'a58' → 58, 'a119' → 119. Returns None if the string cannot be parsed. """ - m = re.search(r"-[a-zA-Z]+(\d+)$", fw_ver.strip()) - if m: - return int(m.group(1)) + if fw_ver: + return int(fw_ver.split('a')[-1].lstrip('-')) + return None - _CCQE_MIN_FW_MINOR = 58 @@ -231,7 +230,7 @@ def _is_all_ionic_support_ccqe() -> bool: return False if len(set(versions)) != 1: return False - + for ver in versions: if not _is_firmware_support_ccqe(ver): return False @@ -246,8 +245,10 @@ def is_ccqe_enabled() -> bool: """Return True if CCQE should be enabled (cached after first call).""" global _ccqe_enabled if _ccqe_enabled is None: - _ccqe_enabled = _lib_has_ionic_ccqe() and _is_all_ionic_support_ccqe() - print(f"Ionic _ccqe_enabled: {_ccqe_enabled}") + lib_support = _lib_has_ionic_ccqe() + nic_support = _is_all_ionic_support_ccqe() + _ccqe_enabled = lib_support and nic_support + print(f"Ionic _ccqe_enabled: {_ccqe_enabled} lib_support {lib_support} nic_support: {nic_support}") return _ccqe_enabled From 30c9c193ad5408e289d3b5ade6881ddba7c2d486 Mon Sep 17 00:00:00 2001 From: qizzhang Date: Mon, 18 May 2026 17:39:41 +0800 Subject: [PATCH 27/33] fix ccqe poll & remove redundant code --- .../ionic/ionic_device_primitives.hpp | 320 +----------------- 1 file changed, 12 insertions(+), 308 deletions(-) diff --git a/include/mori/core/transport/rdma/providers/ionic/ionic_device_primitives.hpp b/include/mori/core/transport/rdma/providers/ionic/ionic_device_primitives.hpp index f00f0ba9..d235687e 100644 --- a/include/mori/core/transport/rdma/providers/ionic/ionic_device_primitives.hpp +++ b/include/mori/core/transport/rdma/providers/ionic/ionic_device_primitives.hpp @@ -465,106 +465,27 @@ inline __device__ void UpdateDbrAndRingDbRecv(void* dbrRecAdd /* Completion Queue */ /* ---------------------------------------------------------------------------------------------- */ #ifdef IONIC_CCQE -template <> -inline __device__ int PollCqOnce(void* cqeAddr, uint32_t cqeNum, - uint32_t consIdx, uint32_t* wqeIdx) { - volatile struct ionic_v1_cqe* cqe = reinterpret_cast(cqeAddr); - uint32_t old, msn = HTOBE32(cqe->send.msg_msn); - - MORI_PRINTF("ABH %s:%d here cons %#x msn %#x\n", __func__, __LINE__, consIdx, msn); - while ((msn - consIdx) & 0x800000) { - old = msn; - msn = HTOBE32(cqe->send.msg_msn); - if (msn != old) { - MORI_PRINTF("ABH %s:%d here cons %#x msn %#x\n", __func__, __LINE__, consIdx, msn); - } - } - MORI_PRINTF("ABH %s:%d here - msn %#x\n", __func__, __LINE__, msn); - - *wqeIdx = msn; - - return 0; -} -#else -template <> -inline __device__ int PollCqOnce(void* cqeAddr, uint32_t cqeNum, - uint32_t consIdx, uint32_t* wqeIdx) { - uint32_t cqeIdx = consIdx & (cqeNum - 1); - char* Addr = reinterpret_cast(cqeAddr) + (cqeIdx * sizeof(struct ionic_v1_cqe)); - struct ionic_v1_cqe* cqe = reinterpret_cast(Addr); - - MORI_PRINTF("ABH %s:%d consIdx:%u, cqeIdx:%u, cqeAddr:%p, qtf_be:0x%08x, cqe->status_length:%d\n", - __func__, __LINE__, consIdx, cqeIdx, Addr, - *(volatile uint32_t*)(&cqe->qid_type_flags), HTOBE32(cqe->status_length)); -#if 1 - MORI_PRINTF("dump cqe at addr:%p\n", Addr); - for (int i = 0; i < 32; i++) { - MORI_PRINTF("%02x", (unsigned char)Addr[i]); - if ((i + 1) % 4 == 0) MORI_PRINTF("\n"); - } -#endif - /* Determine expected color based on cq wrap count */ - uint32_t qtf_color_bit = HTOBE32(IONIC_V1_CQE_COLOR); - uint32_t qtf_color_exp = qtf_color_bit; - if (cqeIdx & cqeNum) { - qtf_color_exp = 0; - } - - /* Check if my cqe color == expected color */ - // first round: 1 == 1, second round: 0 == 0 - uint32_t qtf_be = *(volatile uint32_t*)(&cqe->qid_type_flags); - if ((qtf_be & qtf_color_bit) != qtf_color_exp) { - MORI_PRINTF("cqe not ready\n"); - return -1; // CQE just not ready yet, try again - } +inline __device__ int PollCqCcqe(void* cqAddr, uint32_t cqeNum, uint32_t* consIdx, + uint32_t* wqeCounter) { + const uint32_t curConsIdx = *consIdx; - uint32_t msn = HTOBE32(cqe->send.msg_msn); - - /* Report if the completion indicates an error. */ - if (!!(qtf_be & HTOBE32(IONIC_V1_CQE_ERROR))) { - uint32_t qtf = HTOBE32(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 = cqe->status_length; - uint64_t npg = cqe->send.npg_wqe_idx_timestamp & IONIC_V1_CQE_WQE_IDX_MASK; - MORI_PRINTF("QUIET ERROR: qid %u type %u flag %#x status %u msn %u npg %lu\n", qid, type, flag, - status, msn, npg); - return HTOBE32(cqe->status_length); + volatile struct ionic_v1_cqe* cqe = reinterpret_cast(cqAddr); + const uint32_t msn = BE32TOH(*(volatile uint32_t*)(&cqe->send.msg_msn)); + if ((msn - (curConsIdx + 1)) & 0x800000) { + return -1; // firmware hasn't produced enough completions yet } - - MORI_PRINTF("poll cqe one, success\n"); - + *wqeCounter = msn; return 0; } #endif -template <> -inline __device__ int PollCq(void* cqAddr, uint32_t cqeNum, uint32_t* consIdx) { - const uint32_t curConsIdx = atomicAdd(consIdx, 1); - int err = -1; - - // ABH: polls until each thread sees a ready cqe - // (what if not all threads see a ready cqe?) - do { - err = PollCqOnce(cqAddr, cqeNum, curConsIdx, nullptr); - // TODO: Explain clearly why adding a compiler barrier fix hang issue - asm volatile("" ::: "memory"); - } while (err < 0); - - // Handle error cases - if (err) { - auto error = IonicHandleErrorCqe(err); - MORI_PRINTF("[IONIC PollCq] CQE error: %s (opcode: %d) at %s:%d\n", IbvWcStatusString(error), - err, __FILE__, __LINE__); - return err; - } - - return 0; -} template <> inline __device__ int PollCq(void* cqAddr, uint32_t cqeNum, uint32_t* consIdx, uint32_t* wqeCounter) { +#ifdef IONIC_CCQE + return PollCqCcqe(cqAddr, cqeNum, consIdx, wqeCounter); +#endif + const uint32_t curConsIdx = *consIdx; const uint32_t cqeIdx = curConsIdx & (cqeNum - 1); @@ -591,23 +512,6 @@ inline __device__ int PollCq(void* cqAddr, uint32_t cqeNum, u const uint32_t msn = BE32TOH(cqe->send.msg_msn) & 0xFFFF; const uint8_t error = IonicHandleErrorCqe(status); - // MORI_PRINTF( - // "PollCqOnce2, QUIET ERROR: block:%u, warp:%u, lane:%u, cqeAddr:%p, error:%u " - // "qid %u type %u flag %#x status 0x%08x msn %u npg %lu\n", - // blockIdx.x, threadIdx.x / warpSize, __lane_id(), cqeAddr, error, qid, type, flags, - // status, msn, npg); - -#if 0 - // Debug: dump raw CQE contents - MORI_PRINTF("dump cqe at addr:%p\n", cqeAddr); - for (int i = 0; i < 32; i++) { - MORI_PRINTF("%02x", static_cast(cqeAddr[i])); - if ((i + 1) % 4 == 0) { - MORI_PRINTF("\n"); - } - } -#endif - return error; } @@ -615,193 +519,6 @@ inline __device__ int PollCq(void* cqAddr, uint32_t cqeNum, u return 0; } -#ifdef IONIC_CCQE -inline __device__ int PollCqOnce2(WorkQueueHandle& wqHandle, CompletionQueueHandle& cqHandle, - uint64_t activemask, void* cqeAddr, uint32_t cqeNum, - uint32_t consIdx) { - volatile struct ionic_v1_cqe* cqe = reinterpret_cast(cqeAddr); - uint32_t old, msn = HTOBE32(cqe->send.msg_msn); - - consIdx = wqHandle.dbTouchIdx; - - // MORI_PRINTF("ABH %s:%d here cons %#x msn %#x\n", __func__, __LINE__, consIdx, msn); - while ((msn - consIdx) & 0x800000) { - old = msn; - msn = HTOBE32(cqe->send.msg_msn); - if (msn != old) { - // MORI_PRINTF("ABH %s:%d here cons %#x msn %#x\n", __func__, __LINE__, consIdx, msn); - } - } - - wqHandle.doneIdx = msn; - return 0; -} -#else -inline __device__ int PollCqOnce2(WorkQueueHandle& wqHandle, CompletionQueueHandle& cqHandle, - uint64_t activemask, void* cqeAddr, uint32_t cqeNum, - uint32_t consIdx) { - uint32_t my_logical_lane_id = get_active_lane_num(activemask); - uint32_t my_cq_pos = cqHandle.cq_consumer + my_logical_lane_id; - - uint32_t cqeIdx = my_cq_pos & (cqeNum - 1); - char* Addr = reinterpret_cast(cqeAddr) + (cqeIdx * sizeof(struct ionic_v1_cqe)); - - struct ionic_v1_cqe* cqe = reinterpret_cast(Addr); -#if 0 - MORI_PRINTF("PollCqOnce2, block:%u, warp:%u, lane:%u, consIdx:%u, cqeIdx:%u, cqeAddr:%p, qtf_be:0x%08x, cqe->status_length:%d, msn:%u\n", - blockIdx.x, threadIdx.x/warpSize, __lane_id(), my_cq_pos, cqeIdx, Addr, - *(volatile uint32_t *)(&cqe->qid_type_flags), BE32TOH(cqe->status_length), BE32TOH(cqe->send.msg_msn)); -#endif -#if 0 - MORI_PRINTF("dump cqe at addr:%p\n", Addr); - for (int i = 0; i < 32; i++) { - MORI_PRINTF("%02x", (unsigned char)Addr[i]); - if ((i+1)%4 == 0) - MORI_PRINTF("\n"); - } -#endif - /* Determine expected color based on cq wrap count */ - uint32_t qtf_color_bit = IONIC_V1_CQE_COLOR; - uint32_t qtf_color_exp = qtf_color_bit; - if (my_cq_pos & cqeNum) { - qtf_color_exp = 0; - } - - /* Check if my cqe color == expected color */ - // first round: 1 == 1, second round: 0 == 0 - uint32_t qtf_be = BE32TOH(*(volatile uint32_t*)(&cqe->qid_type_flags)); - if ((qtf_be & qtf_color_bit) != qtf_color_exp) { -#if 0 - MORI_PRINTF("PollCqOnce2, not ready, block:%u, warp:%u, lane:%u, consIdx:%u, cqeIdx:%u, cqeAddr:%p, qtf_be:0x%08x, cqe->status_length:0x%08x, msn:%u\n", - blockIdx.x, threadIdx.x/warpSize, __lane_id(), my_cq_pos, cqeIdx, Addr, - *(volatile uint32_t *)(&cqe->qid_type_flags), BE32TOH(cqe->status_length), BE32TOH(cqe->send.msg_msn)); -#endif - return 0; // CQE just not ready yet, try again - } - - uint32_t msn = BE32TOH(cqe->send.msg_msn); - - /* Report if the completion indicates an error. */ - if (!!(qtf_be & IONIC_V1_CQE_ERROR)) { - uint32_t qtf = 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 = cqe->status_length; - uint64_t npg = cqe->send.npg_wqe_idx_timestamp & IONIC_V1_CQE_WQE_IDX_MASK; - uint8_t error = IonicHandleErrorCqe(BE32TOH(cqe->status_length)); - MORI_PRINTF( - "PollCqOnce2, QUIET ERROR: block:%u, warp:%u, lane:%u, cqeAddr:%p, error:%u qid %u type %u " - "flag %#x status 0x%08x msn %u npg %lu\n", - blockIdx.x, threadIdx.x / warpSize, __lane_id(), Addr, error, qid, type, flag, status, msn, - npg); -#if 1 - MORI_PRINTF("dump cqe at addr:%p\n", Addr); - for (int i = 0; i < 32; i++) { - MORI_PRINTF("%02x", (unsigned char)Addr[i]); - if ((i + 1) % 4 == 0) MORI_PRINTF("\n"); - } -#endif - /* No other way to signal an error, so just crash. */ - // abort(); - return error; - } - -#if 0 - MORI_PRINTF("PollCqOnce2, success, block:%u, warp:%u, lane:%u, qp:%u, cqeAddr:%p, my_cq_pos:%u, cqeNum:%u, msn:%u\n", - blockIdx.x, threadIdx.x/warpSize, __lane_id(), - qtf_be >> IONIC_V1_CQE_QID_SHIFT, Addr, my_cq_pos, cqHandle.cqeNum, msn); -#endif - /* Only proceed with the furthest ahead cqe to update the sq state */ - uint64_t my_lane_mask = 1ull << __lane_id(); - uint64_t lesser_lane_mask = my_lane_mask - 1; - if (my_lane_mask != (__ballot(true) & activemask & ~lesser_lane_mask)) { - return 0; - } - - /* update position in the cq */ - cqHandle.cq_consumer = my_cq_pos + 1; - - /* - * Ring cq doorbell frequently enough to avoid cq full. - * - * NB: IONIC_CQ_GRACE is 100 - */ - if (((cqHandle.cq_consumer - cqHandle.cq_dbpos) & (cqHandle.cqeNum - 1)) >= 100) { - cqHandle.cq_dbpos = cqHandle.cq_consumer; - uint64_t dbrVal = cqHandle.cq_dbval | ((cqHandle.cqeNum - 1) & (cqHandle.cq_dbpos)); -#if 0 - MORI_PRINTF("update cq doorbell, block:%u, warp:%u, lane:%u, cq dbrAddr:%p, dbrVal:0x%lx, cq_consumer:%u\n", - blockIdx.x, threadIdx.x/warpSize, __lane_id(), reinterpret_cast(cqHandle.dbrRecAddr), dbrVal, cqHandle.cq_consumer); -#endif - __atomic_store_n(reinterpret_cast(cqHandle.dbrRecAddr), dbrVal, - __ATOMIC_SEQ_CST); // TODO:maybe relaxed? - } - - wqHandle.doneIdx = msn; - return 0; -} -#endif - -#ifdef IONIC_CCQE -template <> -inline __device__ int PollCq(WorkQueueHandle& wqHandle, - CompletionQueueHandle& cqHandle, void* cqAddr, - uint32_t cqeNum, uint32_t* consIdx, - uint16_t* wqeCounter) { - PollCqOnce2(wqHandle, cqHandle, 1, cqAddr, cqeNum, *consIdx); - *wqeCounter = *consIdx; - return 0; -} -#else -template <> -inline __device__ int PollCq(WorkQueueHandle& wqHandle, - CompletionQueueHandle& cqHandle, void* cqAddr, - uint32_t cqeNum, uint32_t* consIdx, - uint16_t* wqeCounter) { - uint32_t greed = 10; - const uint32_t curConsIdx = *consIdx; - uint64_t activemask = GetActiveLaneMask(); - uint32_t cons = wqHandle.dbTouchIdx; - int err; - /* wait for sq_msn to catch up or pass cons. */ - /* 0x800000 - sign bit for 24-bit fields */ - while ((wqHandle.doneIdx - cons) & 0x800000) { - if (!spin_lock_try_acquire_shared(&cqHandle.pollCqLock, activemask)) { - continue; - } - - /* with lock acquired, this wave polls cqes until caught up */ - while ((wqHandle.doneIdx - cons) & 0x800000) { - uint32_t old_sq_msn = wqHandle.doneIdx; - // MORI_PRINTF("PollCq, before PollCqOnce2, curConsIdx:%u\n", curConsIdx); - // asm volatile("" ::: "memory"); - err = PollCqOnce2(wqHandle, cqHandle, activemask, cqAddr, cqeNum, curConsIdx); - if (err != 0) { - MORI_PRINTF("PollCq, PollCqOnce2 failed, err:%u\n", err); - return err; - } - asm volatile("" ::: "memory"); - // MORI_PRINTF("PollCq, after PollCqOnce2, curConsIdx:%u\n", curConsIdx); - if (!((wqHandle.doneIdx - cons) & 0x800000)) { - if (wqHandle.doneIdx == old_sq_msn) { - break; - } - if (!greed) { - break; - } - --greed; - } - } - - spin_lock_release_shared(&cqHandle.pollCqLock, activemask); - break; - } - - return 0; -} -#endif - template <> inline __device__ void UpdateCqDbrRecord(CompletionQueueHandle& cq, uint32_t consIdx) { @@ -814,19 +531,6 @@ inline __device__ void UpdateCqDbrRecord(CompletionQueueHandl #endif } -template <> -inline __device__ int PollCqAndUpdateDbr(CompletionQueueHandle& cq, - uint32_t* consIdx, uint32_t* lockVar) { - AcquireLock(lockVar); - - int err = PollCq(cq.cqAddr, cq.cqeNum, consIdx); - if (err >= 0) { - UpdateCqDbrRecord(cq, *consIdx); - } - - ReleaseLock(lockVar); - return err; -} // #endif } // namespace core } // namespace mori From 64fe67097f84997bb1e31b55e1141ebb4aeebc11 Mon Sep 17 00:00:00 2001 From: Qizhou Zhang Date: Mon, 18 May 2026 12:54:20 +0000 Subject: [PATCH 28/33] fix local rdma example --- examples/local_rdma_ops/atomic_gpu.cpp | 14 +++++++---- examples/local_rdma_ops/send_recv_gpu.cpp | 25 +++++++++++++++----- examples/local_rdma_ops/write_inline_gpu.cpp | 14 ++++++++--- 3 files changed, 40 insertions(+), 13 deletions(-) diff --git a/examples/local_rdma_ops/atomic_gpu.cpp b/examples/local_rdma_ops/atomic_gpu.cpp index 1ccdeea2..2cf13096 100644 --- a/examples/local_rdma_ops/atomic_gpu.cpp +++ b/examples/local_rdma_ops/atomic_gpu.cpp @@ -72,18 +72,24 @@ __device__ void SendThreadKernel(RdmaEndpoint& epSend, RdmaMemoryRegion sendMr, } __device__ void RecvThreadKernel(RdmaEndpoint& epRecv, RdmaMemoryRegion mr) { - uint32_t postIdx = 0; uint32_t* addr = reinterpret_cast(mr.addr); uint32_t val = core::AtomicLoadSeqCst(addr); printf("val = %u\n", val); - while (val != 2) { + + // Cross-block, lock-free observation: there is no sync between the send block + // (which issues CAS then FETCH_ADD) and this recv block, so by the time we + // start polling either both atomics or only CAS may have landed. Just wait + // until the value leaves its initial 0, and then until it reaches the final + // expected sum (CAS_swap + FETCH_ADD_value == 2 + 2 == 4). + while (val == 0) { val = core::AtomicLoadSeqCst(addr); - printf("after compare and swap val = %u\n", val); } + printf("after compare and swap val = %u\n", val); + while (val != 4) { val = core::AtomicLoadSeqCst(addr); - printf("after fetch add val = %u\n", val); } + printf("after fetch add val = %u\n", val); } __global__ void SendRecvOnGpu(RdmaEndpoint& epSend, RdmaEndpoint& epRecv, RdmaMemoryRegion mrSend, diff --git a/examples/local_rdma_ops/send_recv_gpu.cpp b/examples/local_rdma_ops/send_recv_gpu.cpp index a880fe52..12c80264 100644 --- a/examples/local_rdma_ops/send_recv_gpu.cpp +++ b/examples/local_rdma_ops/send_recv_gpu.cpp @@ -53,9 +53,17 @@ __device__ void SendThreadKernel(RdmaEndpoint& epSend, RdmaMemoryRegion mr, int printf("RingDoorbell is done\n"); __threadfence_system(); - int snd_opcode = - PollCq

(epSend.cqHandle.cqAddr, epSend.cqHandle.cqeNum, &epSend.cqHandle.consIdx); - printf("send PollCq is done\n"); + // PSD 4-arg PollCq is non-blocking (returns -1 when the CCQE msg_msn isn't + // there yet); BNXT/MLX5 spin internally. Wrap in a busy-wait loop so this + // example works uniformly across all providers. + uint32_t snd_wqeIdx = 0; + int snd_opcode; + do { + snd_opcode = PollCq

(epSend.cqHandle.cqAddr, epSend.cqHandle.cqeNum, + &epSend.cqHandle.consIdx, &snd_wqeIdx); + } while (snd_opcode < 0); + epSend.cqHandle.consIdx += 1; + printf("send PollCq is done, wqeIdx %u\n", snd_wqeIdx); UpdateCqDbrRecord

(epSend.cqHandle, epSend.cqHandle.consIdx); printf("send UpdateCqDbrRecord is done\n"); // printf("snd_opcode %d val %d\n", snd_opcode, reinterpret_cast(mrSend.addr)[0]); @@ -84,9 +92,14 @@ __device__ void RecvThreadKernel(RdmaEndpoint& epRecv, RdmaMemoryRegion mr, int printf("recv RingDoorbell is done\n"); } - int rcv_opcode = - PollCq

(epRecv.cqHandle.cqAddr, epRecv.cqHandle.cqeNum, &epRecv.cqHandle.consIdx); - printf("recv PollCq is done\n"); + uint32_t rcv_wqeIdx = 0; + int rcv_opcode; + do { + rcv_opcode = PollCq

(epRecv.cqHandle.cqAddr, epRecv.cqHandle.cqeNum, + &epRecv.cqHandle.consIdx, &rcv_wqeIdx); + } while (rcv_opcode < 0); + epRecv.cqHandle.consIdx += 1; + printf("recv PollCq is done, wqeIdx %u\n", rcv_wqeIdx); UpdateCqDbrRecord

(epRecv.cqHandle, epRecv.cqHandle.consIdx); printf("recv UpdateCqDbrRecord is done\n"); diff --git a/examples/local_rdma_ops/write_inline_gpu.cpp b/examples/local_rdma_ops/write_inline_gpu.cpp index f54cc7a4..a3e0c20b 100644 --- a/examples/local_rdma_ops/write_inline_gpu.cpp +++ b/examples/local_rdma_ops/write_inline_gpu.cpp @@ -53,11 +53,19 @@ __device__ void SendThreadKernel(RdmaEndpoint& epSend, RdmaMemoryRegion mr) { RingDoorbell

(epSend.wqHandle.dbrAddr, dbr_val); __threadfence_system(); - int opcode = - PollCq

(epSend.cqHandle.cqAddr, epSend.cqHandle.cqeNum, &epSend.cqHandle.consIdx); + // PSD 4-arg PollCq is non-blocking (returns -1 when CQE / CCQE msg_msn not + // ready yet), while BNXT/MLX5 internally spin. Wrap in a busy-wait loop so + // this example works uniformly across all providers. + uint32_t wqeIdx = 0; + int opcode; + do { + opcode = PollCq

(epSend.cqHandle.cqAddr, epSend.cqHandle.cqeNum, + &epSend.cqHandle.consIdx, &wqeIdx); + } while (opcode < 0); + epSend.cqHandle.consIdx += 1; __threadfence_system(); UpdateCqDbrRecord

(epSend.cqHandle, epSend.cqHandle.consIdx); - // printf("round %d snd_opcode %d\n", i, opcode); + // printf("round %d snd_opcode %d wqeIdx %u\n", i, opcode, wqeIdx); raddr += i; } From 603925c7ce858278f3297e5f05fa1759668a7f20 Mon Sep 17 00:00:00 2001 From: qizzhang Date: Tue, 19 May 2026 10:17:08 +0800 Subject: [PATCH 29/33] format code --- examples/local_rdma_ops/write_inline_gpu.cpp | 4 ++-- python/mori/jit/core.py | 11 +++++++---- 2 files changed, 9 insertions(+), 6 deletions(-) diff --git a/examples/local_rdma_ops/write_inline_gpu.cpp b/examples/local_rdma_ops/write_inline_gpu.cpp index a3e0c20b..bcdd1cbe 100644 --- a/examples/local_rdma_ops/write_inline_gpu.cpp +++ b/examples/local_rdma_ops/write_inline_gpu.cpp @@ -59,8 +59,8 @@ __device__ void SendThreadKernel(RdmaEndpoint& epSend, RdmaMemoryRegion mr) { uint32_t wqeIdx = 0; int opcode; do { - opcode = PollCq

(epSend.cqHandle.cqAddr, epSend.cqHandle.cqeNum, - &epSend.cqHandle.consIdx, &wqeIdx); + opcode = PollCq

(epSend.cqHandle.cqAddr, epSend.cqHandle.cqeNum, &epSend.cqHandle.consIdx, + &wqeIdx); } while (opcode < 0); epSend.cqHandle.consIdx += 1; __threadfence_system(); diff --git a/python/mori/jit/core.py b/python/mori/jit/core.py index 8062b1bd..369c4d9b 100644 --- a/python/mori/jit/core.py +++ b/python/mori/jit/core.py @@ -185,10 +185,11 @@ def _parse_ionic_fw_minor(fw_ver: str) -> int | None: Returns None if the string cannot be parsed. """ if fw_ver: - return int(fw_ver.split('a')[-1].lstrip('-')) - + return int(fw_ver.split("a")[-1].lstrip("-")) + return None + _CCQE_MIN_FW_MINOR = 58 @@ -230,7 +231,7 @@ def _is_all_ionic_support_ccqe() -> bool: return False if len(set(versions)) != 1: return False - + for ver in versions: if not _is_firmware_support_ccqe(ver): return False @@ -248,7 +249,9 @@ def is_ccqe_enabled() -> bool: lib_support = _lib_has_ionic_ccqe() nic_support = _is_all_ionic_support_ccqe() _ccqe_enabled = lib_support and nic_support - print(f"Ionic _ccqe_enabled: {_ccqe_enabled} lib_support {lib_support} nic_support: {nic_support}") + print( + f"Ionic _ccqe_enabled: {_ccqe_enabled} lib_support {lib_support} nic_support: {nic_support}" + ) return _ccqe_enabled From 3bb0ec67965913b6dc4241f6d78b3d35b73d153b Mon Sep 17 00:00:00 2001 From: qizzhang Date: Wed, 20 May 2026 10:04:25 +0800 Subject: [PATCH 30/33] use env var to control ccqe mode --- python/mori/jit/core.py | 16 ++++++++++------ .../transport/rdma/providers/ionic/ionic.cpp | 2 ++ 2 files changed, 12 insertions(+), 6 deletions(-) diff --git a/python/mori/jit/core.py b/python/mori/jit/core.py index 369c4d9b..cfc60cc6 100644 --- a/python/mori/jit/core.py +++ b/python/mori/jit/core.py @@ -246,12 +246,16 @@ def is_ccqe_enabled() -> bool: """Return True if CCQE should be enabled (cached after first call).""" global _ccqe_enabled if _ccqe_enabled is None: - lib_support = _lib_has_ionic_ccqe() - nic_support = _is_all_ionic_support_ccqe() - _ccqe_enabled = lib_support and nic_support - print( - f"Ionic _ccqe_enabled: {_ccqe_enabled} lib_support {lib_support} nic_support: {nic_support}" - ) + if os.environ.get("MORI_DISABLE_IONIC_CCQE"): + _ccqe_enabled = False + print("Ionic _ccqe_enabled: False (disabled by MORI_DISABLE_IONIC_CCQE)") + else: + lib_support = _lib_has_ionic_ccqe() + nic_support = _is_all_ionic_support_ccqe() + _ccqe_enabled = lib_support and nic_support + print( + f"Ionic _ccqe_enabled: {_ccqe_enabled} lib_support {lib_support} nic_support: {nic_support}" + ) return _ccqe_enabled diff --git a/src/application/transport/rdma/providers/ionic/ionic.cpp b/src/application/transport/rdma/providers/ionic/ionic.cpp index f803bb69..d53453a2 100644 --- a/src/application/transport/rdma/providers/ionic/ionic.cpp +++ b/src/application/transport/rdma/providers/ionic/ionic.cpp @@ -27,6 +27,7 @@ #include #include +#include #include #include @@ -70,6 +71,7 @@ int ReadIonicFwBuild(const char* dev_name) { } bool IsCcqeSupported(ibv_context* context) { + if (std::getenv("MORI_DISABLE_IONIC_CCQE")) return false; if (IonicDvApi::Instance().create_cq_ex == nullptr) return false; int build = ReadIonicFwBuild(context->device->name); From 9f23058e86c04656bf154bd5078e27aebfd81a89 Mon Sep 17 00:00:00 2001 From: qizzhang Date: Wed, 20 May 2026 10:44:11 +0800 Subject: [PATCH 31/33] judge env value --- python/mori/jit/core.py | 2 +- src/application/transport/rdma/providers/ionic/ionic.cpp | 4 +++- 2 files changed, 4 insertions(+), 2 deletions(-) diff --git a/python/mori/jit/core.py b/python/mori/jit/core.py index cfc60cc6..8a005cbf 100644 --- a/python/mori/jit/core.py +++ b/python/mori/jit/core.py @@ -246,7 +246,7 @@ def is_ccqe_enabled() -> bool: """Return True if CCQE should be enabled (cached after first call).""" global _ccqe_enabled if _ccqe_enabled is None: - if os.environ.get("MORI_DISABLE_IONIC_CCQE"): + if os.environ.get("MORI_DISABLE_IONIC_CCQE") == "1": _ccqe_enabled = False print("Ionic _ccqe_enabled: False (disabled by MORI_DISABLE_IONIC_CCQE)") else: diff --git a/src/application/transport/rdma/providers/ionic/ionic.cpp b/src/application/transport/rdma/providers/ionic/ionic.cpp index d53453a2..3d2abb62 100644 --- a/src/application/transport/rdma/providers/ionic/ionic.cpp +++ b/src/application/transport/rdma/providers/ionic/ionic.cpp @@ -28,6 +28,7 @@ #include #include #include +#include #include #include @@ -71,7 +72,8 @@ int ReadIonicFwBuild(const char* dev_name) { } bool IsCcqeSupported(ibv_context* context) { - if (std::getenv("MORI_DISABLE_IONIC_CCQE")) return false; + const char* disable_ccqe = std::getenv("MORI_DISABLE_IONIC_CCQE"); + if (disable_ccqe && std::strcmp(disable_ccqe, "1") == 0) return false; if (IonicDvApi::Instance().create_cq_ex == nullptr) return false; int build = ReadIonicFwBuild(context->device->name); From ddea7f22f1fb4b8d66c16716e23598b525ca3a9a Mon Sep 17 00:00:00 2001 From: qizzhang Date: Thu, 21 May 2026 11:28:17 +0800 Subject: [PATCH 32/33] address jhaos comments --- .../ionic/ionic_device_primitives.hpp | 14 +++---- .../transport/rdma/providers/ionic/ionic_fw.h | 2 +- python/mori/jit/core.py | 30 +++++++------- .../transport/rdma/providers/ionic/ionic.cpp | 41 +++++++++++-------- 4 files changed, 48 insertions(+), 39 deletions(-) diff --git a/include/mori/core/transport/rdma/providers/ionic/ionic_device_primitives.hpp b/include/mori/core/transport/rdma/providers/ionic/ionic_device_primitives.hpp index d235687e..fd68159e 100644 --- a/include/mori/core/transport/rdma/providers/ionic/ionic_device_primitives.hpp +++ b/include/mori/core/transport/rdma/providers/ionic/ionic_device_primitives.hpp @@ -465,8 +465,9 @@ inline __device__ void UpdateDbrAndRingDbRecv(void* dbrRecAdd /* Completion Queue */ /* ---------------------------------------------------------------------------------------------- */ #ifdef IONIC_CCQE -inline __device__ int PollCqCcqe(void* cqAddr, uint32_t cqeNum, uint32_t* consIdx, - uint32_t* wqeCounter) { +template <> +inline __device__ int PollCq(void* cqAddr, uint32_t cqeNum, uint32_t* consIdx, + uint32_t* wqeCounter) { const uint32_t curConsIdx = *consIdx; volatile struct ionic_v1_cqe* cqe = reinterpret_cast(cqAddr); @@ -474,18 +475,16 @@ inline __device__ int PollCqCcqe(void* cqAddr, uint32_t cqeNum, uint32_t* consId if ((msn - (curConsIdx + 1)) & 0x800000) { return -1; // firmware hasn't produced enough completions yet } + *wqeCounter = msn; return 0; } -#endif + +#else template <> inline __device__ int PollCq(void* cqAddr, uint32_t cqeNum, uint32_t* consIdx, uint32_t* wqeCounter) { -#ifdef IONIC_CCQE - return PollCqCcqe(cqAddr, cqeNum, consIdx, wqeCounter); -#endif - const uint32_t curConsIdx = *consIdx; const uint32_t cqeIdx = curConsIdx & (cqeNum - 1); @@ -518,6 +517,7 @@ inline __device__ int PollCq(void* cqAddr, uint32_t cqeNum, u *wqeCounter = BE32TOH(cqe->send.msg_msn); return 0; } +#endif // end of IONIC_CCQE template <> inline __device__ void UpdateCqDbrRecord(CompletionQueueHandle& cq, diff --git a/include/mori/core/transport/rdma/providers/ionic/ionic_fw.h b/include/mori/core/transport/rdma/providers/ionic/ionic_fw.h index 51679f7a..7bc40919 100644 --- a/include/mori/core/transport/rdma/providers/ionic/ionic_fw.h +++ b/include/mori/core/transport/rdma/providers/ionic/ionic_fw.h @@ -243,7 +243,7 @@ enum ionic_v1_cqe_qtf_bits { IONIC_V1_CQE_TYPE_RECV = 1, IONIC_V1_CQE_TYPE_SEND_MSN = 2, IONIC_V1_CQE_TYPE_SEND_NPG = 3, - c = 4, + IONIC_V1_CQE_TYPE_RECV_RCQE = 4, }; #if !defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_HCC__) diff --git a/python/mori/jit/core.py b/python/mori/jit/core.py index 8a005cbf..5b4c1688 100644 --- a/python/mori/jit/core.py +++ b/python/mori/jit/core.py @@ -178,25 +178,23 @@ def _lib_has_ionic_ccqe() -> bool: return False -def _parse_ionic_fw_minor(fw_ver: str) -> int | None: - """Parse the build number from an ionic firmware string like '1.117.5-a58'. - - Extracts the numeric part from the suffix after '-', e.g. 'a58' → 58, 'a119' → 119. - Returns None if the string cannot be parsed. - """ - if fw_ver: - return int(fw_ver.split("a")[-1].lstrip("-")) - - return None +_CCQE_MIN_FW_VERSION = (1, 117, 5, 58) -_CCQE_MIN_FW_MINOR = 58 +def _parse_ionic_fw_version(fw_ver: str) -> tuple[int, ...] | None: + """Parse '1.117.5-a-58' → (1, 117, 5, 58). Returns None if unparseable.""" + if not fw_ver: + return None + m = re.match(r"^(\d+)\.(\d+)\.(\d+)-a-?(\d+)$", fw_ver) + if not m: + return None + return tuple(int(x) for x in m.groups()) def _is_firmware_support_ccqe(fw_ver: str) -> bool: - """Return True if the firmware version string reports build number >= 58.""" - minor = _parse_ionic_fw_minor(fw_ver) - return minor is not None and minor >= _CCQE_MIN_FW_MINOR + """Return True if the firmware version >= 1.117.5-a-58.""" + ver = _parse_ionic_fw_version(fw_ver) + return ver is not None and ver >= _CCQE_MIN_FW_VERSION def _get_ionic_fw_versions() -> list[str]: @@ -232,6 +230,8 @@ def _is_all_ionic_support_ccqe() -> bool: if len(set(versions)) != 1: return False + print(f"ionic ver: {versions[-1]}") + for ver in versions: if not _is_firmware_support_ccqe(ver): return False @@ -410,7 +410,7 @@ def _hipcc_genco( *_ccqe_defines(), *_profiler_defines(), ] - print("genco cmd", cmd) + for d in include_dirs: cmd.extend(["-I", str(d)]) cmd.extend([str(source), "-o", str(output)]) diff --git a/src/application/transport/rdma/providers/ionic/ionic.cpp b/src/application/transport/rdma/providers/ionic/ionic.cpp index 3d2abb62..d0a2c4cc 100644 --- a/src/application/transport/rdma/providers/ionic/ionic.cpp +++ b/src/application/transport/rdma/providers/ionic/ionic.cpp @@ -31,6 +31,7 @@ #include #include #include +#include #include "mori/application/utils/check.hpp" #include "mori/application/utils/math.hpp" @@ -46,39 +47,47 @@ namespace application { namespace { -// Minimum firmware build number required for CCQE support (e.g. "1.117.5-a58" → 58). -constexpr int kCcqeMinFwBuild = 58; +using FwVersion = std::tuple; +constexpr FwVersion kCcqeMinFwVersion{1, 117, 5, 58}; -// Read /sys/class/infiniband//fw_ver and return the numeric build suffix -// (the digits after the last '-[letters]' component). Returns -1 on any failure. -int ReadIonicFwBuild(const char* dev_name) { +// Parse "1.117.5-a-58" or "1.117.5-a58" into (1,117,5,58). +std::optional ParseIonicFwVersion(const char* fw_ver) { + int major, minor, patch, build; + char tag; + if (sscanf(fw_ver, "%d.%d.%d-%c-%d", &major, &minor, &patch, &tag, &build) == 5 || + sscanf(fw_ver, "%d.%d.%d-%c%d", &major, &minor, &patch, &tag, &build) == 5) { + return FwVersion{major, minor, patch, build}; + } + return std::nullopt; +} + +std::optional ReadIonicFwVersion(const char* dev_name) { char path[256]; snprintf(path, sizeof(path), "/sys/class/infiniband/%s/fw_ver", dev_name); FILE* f = fopen(path, "r"); - if (!f) return -1; + if (!f) return std::nullopt; char buf[64] = {}; fgets(buf, sizeof(buf), f); fclose(f); - // Find last '-' then skip letters to reach the build digits. - char* dash = strrchr(buf, '-'); - if (!dash) return -1; - char* p = dash + 1; - while (*p && !isdigit(static_cast(*p))) ++p; - if (!*p) return -1; - return atoi(p); + // Strip trailing newline. + buf[strcspn(buf, "\n")] = '\0'; + return ParseIonicFwVersion(buf); } bool IsCcqeSupported(ibv_context* context) { const char* disable_ccqe = std::getenv("MORI_DISABLE_IONIC_CCQE"); if (disable_ccqe && std::strcmp(disable_ccqe, "1") == 0) return false; if (IonicDvApi::Instance().create_cq_ex == nullptr) return false; - int build = ReadIonicFwBuild(context->device->name); - MORI_APP_TRACE("dev: {} fw_build {}", context->device->name, build); - return build >= kCcqeMinFwBuild; + /* Minimum firmware version verified by MORI to support CCQE is 1.117.5-a-58. */ + auto ver = ReadIonicFwVersion(context->device->name); + MORI_APP_TRACE("dev: {} fw_ver {}.{}.{}-a-{}", context->device->name, + ver ? std::get<0>(*ver) : -1, ver ? std::get<1>(*ver) : -1, + ver ? std::get<2>(*ver) : -1, ver ? std::get<3>(*ver) : -1); + return ver.has_value() && *ver >= kCcqeMinFwVersion; } } // namespace From 7edc3b1880943e23608d8177918b60bc1841a5b0 Mon Sep 17 00:00:00 2001 From: qizzhang Date: Thu, 21 May 2026 14:55:42 +0800 Subject: [PATCH 33/33] remove some code --- .../mori/core/transport/rdma/providers/ionic/ionic_defs.hpp | 3 +-- .../transport/rdma/providers/ionic/ionic_device_primitives.hpp | 2 +- 2 files changed, 2 insertions(+), 3 deletions(-) diff --git a/include/mori/core/transport/rdma/providers/ionic/ionic_defs.hpp b/include/mori/core/transport/rdma/providers/ionic/ionic_defs.hpp index 08b22fa5..9d3cd9be 100644 --- a/include/mori/core/transport/rdma/providers/ionic/ionic_defs.hpp +++ b/include/mori/core/transport/rdma/providers/ionic/ionic_defs.hpp @@ -24,9 +24,8 @@ namespace mori { namespace core { -// Remove CCQE Defines #define QUEUE_SIZE 1 #define MAX_INLINE_SIZE 32 -// #define IONIC_CCQE 1 + } // namespace core } // namespace mori diff --git a/include/mori/core/transport/rdma/providers/ionic/ionic_device_primitives.hpp b/include/mori/core/transport/rdma/providers/ionic/ionic_device_primitives.hpp index fd68159e..f4397412 100644 --- a/include/mori/core/transport/rdma/providers/ionic/ionic_device_primitives.hpp +++ b/include/mori/core/transport/rdma/providers/ionic/ionic_device_primitives.hpp @@ -517,7 +517,7 @@ inline __device__ int PollCq(void* cqAddr, uint32_t cqeNum, u *wqeCounter = BE32TOH(cqe->send.msg_msn); return 0; } -#endif // end of IONIC_CCQE +#endif // end of PollCq template <> inline __device__ void UpdateCqDbrRecord(CompletionQueueHandle& cq,