[GH-ISSUE #308] ibgda_poll_cq get cqe method #218

New issue

Closed

opened 2026-10-26 09:56:49 +00:00 by bsakaguchi · 1 comment

bsakaguchi commented

2026-10-26 09:56:49 +00:00

Owner

Originally created by @nope8 on GitHub (Jul 16, 2025).
Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/308

Hello,

I can't figure out why always get CQE from cq->cqe which is cq buffer head, not get CQE by cq->cqe + cons_idx.
Does anyone can help me understand this? Thank you very much!

https://github.com/deepseek-ai/DeepEP/blob/main/csrc/kernels/ibgda_device.cuh:472

    // This is a simplified version of NVSHMEM's `ibgda_poll_cq`. 
    // Note that this implementation does not guarantee thread safety,
    // so we must ensure that no other threads are concurrently using the same QP.
    __device__ static __forceinline__ void
    ibgda_poll_cq(nvshmemi_ibgda_device_cq_t *cq, uint64_t idx) {
        const auto cqe64 = static_cast<mlx5_cqe64*>(cq->cqe);
        const uint32_t ncqes = cq->ncqes;
        memory_fence_cta();
    
        // NOTES: this while loop is part of do-while below.
        // `wqe_counter` is the HW consumer index. However, we always maintain `index + 1`.
        // To be able to compare with the index, we need to use `wqe_counter + 1`.
        // Because `wqe_counter` is `uint16_t`, it may be overflow. Still, we know for
        // sure that if `idx - wqe_counter - 1 < ncqes`, `wqe_counter + 1 is less than
        // idx, and thus we need to wait. We don't need to wait when `idx == wqe_counter + 1`
        // That's why we use `- 2` here to make this case overflow.
        uint16_t wqe_counter;
        do {
            wqe_counter = HtoBE16(ld_na_relaxed(&cqe64->wqe_counter));
        } while ((static_cast<uint16_t>(static_cast<uint16_t>(idx) - wqe_counter - static_cast<uint16_t>(2)) < ncqes));
        *cq->cons_idx = idx;
    
        // Prevent reordering of this function and later instructions
        memory_fence_cta();
    }

Originally created by @nope8 on GitHub (Jul 16, 2025). Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/308 Hello, I can't figure out why always get CQE from cq->cqe which is cq buffer head, not get CQE by cq->cqe + cons_idx. Does anyone can help me understand this? Thank you very much! https://github.com/deepseek-ai/DeepEP/blob/main/csrc/kernels/ibgda_device.cuh:472 // This is a simplified version of NVSHMEM's `ibgda_poll_cq`. // Note that this implementation does not guarantee thread safety, // so we must ensure that no other threads are concurrently using the same QP. __device__ static __forceinline__ void ibgda_poll_cq(nvshmemi_ibgda_device_cq_t *cq, uint64_t idx) { const auto cqe64 = static_cast<mlx5_cqe64*>(cq->cqe); const uint32_t ncqes = cq->ncqes; memory_fence_cta(); // NOTES: this while loop is part of do-while below. // `wqe_counter` is the HW consumer index. However, we always maintain `index + 1`. // To be able to compare with the index, we need to use `wqe_counter + 1`. // Because `wqe_counter` is `uint16_t`, it may be overflow. Still, we know for // sure that if `idx - wqe_counter - 1 < ncqes`, `wqe_counter + 1 is less than // idx, and thus we need to wait. We don't need to wait when `idx == wqe_counter + 1` // That's why we use `- 2` here to make this case overflow. uint16_t wqe_counter; do { wqe_counter = HtoBE16(ld_na_relaxed(&cqe64->wqe_counter)); } while ((static_cast<uint16_t>(static_cast<uint16_t>(idx) - wqe_counter - static_cast<uint16_t>(2)) < ncqes)); *cq->cons_idx = idx; // Prevent reordering of this function and later instructions memory_fence_cta(); }

bsakaguchi closed this issue

2026-10-26 09:56:49 +00:00

bsakaguchi commented

2026-10-26 09:56:50 +00:00

Author

Owner

@nope8 commented on GitHub (Jul 16, 2025):

I find the answer， because the cq is created with attribute collapsed flag when call ibgda_create_cq function in ibgda.cpp , which means all CQE are writen collapsed to first CQ entry.
nvshmem_src\src\modules\transport\ibgda\ibgda.cpp:
int ibgda_create_cq(struct ibgda_cq **pgcq, struct ibgda_device *device)
{

**DEVX_SET(cqc, cq_context, cc, 0x1);  // Use collapsed CQ**

}

@nope8 commented on GitHub (Jul 16, 2025): I find the answer， because the cq is created with attribute collapsed flag when call ibgda_create_cq function in ibgda.cpp , which means all CQE are writen collapsed to first CQ entry. nvshmem_src\src\modules\transport\ibgda\ibgda.cpp: int ibgda_create_cq(struct ibgda_cq **pgcq, struct ibgda_device *device) { **DEVX_SET(cqc, cq_context, cc, 0x1); // Use collapsed CQ** } <img width="915" height="424" alt="Image" src="https://github.com/user-attachments/assets/ad30f69e-e22c-4494-9593-c169cf4f2d2b" />

bsakaguchi referenced this issue

2026-10-26 10:01:10 +00:00

[PR #218] Support other NVLink scenarios #474