[GH-ISSUE #70] rq seems to have not been initialized correctly #65

New issue

Closed

opened 2026-10-26 09:53:37 +00:00 by bsakaguchi · 1 comment

bsakaguchi commented

2026-10-26 09:53:37 +00:00

Owner

Originally created by @Thunderbrook on GitHub (Mar 12, 2025).
Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/70

Hi
ibgda_write_empty_recv_wqe try to clear the rq buf with terminate sge

__global__ void ibgda_initialize_recv_queue(int rank) {
    auto thread_idx = static_cast<int>(threadIdx.x);
    auto num_threads = static_cast<int>(blockDim.x);

    auto dst_rank = static_cast<int>(blockIdx.x);
    if (dst_rank != rank) {
        for (int qp_id = thread_idx; qp_id < ibgda_get_state()->num_rc_per_pe; qp_id += num_threads) {
            auto qp = ibgda_get_rc(dst_rank, qp_id);

            // Clean some necessary variables
            for (int i = 0; i < qp->rx_wq.nwqes; ++ i)
                ibgda_write_empty_recv_wqe(ibgda_get_wqe_ptr(qp, i));
            qp->mvars.rx_wq.resv_head = 0;
            qp->mvars.rx_wq.cons_idx = 0;

            // Allocate receive slots
            nvshmemi_ibgda_allocate_recvs(qp);
        }
    }
}

but in ibgda_get_wqe_ptr will get the sq buf

__device__ static __forceinline__ void*
ibgda_get_wqe_ptr(nvshmemi_ibgda_device_qp_t* qp, uint16_t wqe_idx) {
    uint16_t cnt = qp->tx_wq.nwqes;
    uint16_t idx = wqe_idx & (cnt - 1);
    return reinterpret_cast<void*>(reinterpret_cast<uintptr_t>(qp->tx_wq.wqe) + (idx << MLX5_SEND_WQE_SHIFT));
}

thanks

Originally created by @Thunderbrook on GitHub (Mar 12, 2025). Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/70 Hi ibgda_write_empty_recv_wqe try to clear the rq buf with terminate sge ```cpp __global__ void ibgda_initialize_recv_queue(int rank) { auto thread_idx = static_cast<int>(threadIdx.x); auto num_threads = static_cast<int>(blockDim.x); auto dst_rank = static_cast<int>(blockIdx.x); if (dst_rank != rank) { for (int qp_id = thread_idx; qp_id < ibgda_get_state()->num_rc_per_pe; qp_id += num_threads) { auto qp = ibgda_get_rc(dst_rank, qp_id); // Clean some necessary variables for (int i = 0; i < qp->rx_wq.nwqes; ++ i) ibgda_write_empty_recv_wqe(ibgda_get_wqe_ptr(qp, i)); qp->mvars.rx_wq.resv_head = 0; qp->mvars.rx_wq.cons_idx = 0; // Allocate receive slots nvshmemi_ibgda_allocate_recvs(qp); } } } ``` but in ibgda_get_wqe_ptr will get the sq buf ```cpp __device__ static __forceinline__ void* ibgda_get_wqe_ptr(nvshmemi_ibgda_device_qp_t* qp, uint16_t wqe_idx) { uint16_t cnt = qp->tx_wq.nwqes; uint16_t idx = wqe_idx & (cnt - 1); return reinterpret_cast<void*>(reinterpret_cast<uintptr_t>(qp->tx_wq.wqe) + (idx << MLX5_SEND_WQE_SHIFT)); } ``` thanks

bsakaguchi closed this issue

2026-10-26 09:53:37 +00:00

bsakaguchi commented

2026-10-26 09:53:38 +00:00

Author

Owner

@sphish commented on GitHub (Mar 12, 2025):

You're right, thank you for pointing out this error! We are about to update to a version that doesn't require a recv queue.

@sphish commented on GitHub (Mar 12, 2025): You're right, thank you for pointing out this error! We are about to update to a version that doesn't require a recv queue.

No labels

pull-request

No milestone

No project

No assignees

1 participant

Notifications

Due date

The due date is invalid or out of range. Please use the format "yyyy-mm-dd".

No due date set.

Dependencies

No dependencies set.

Reference

deepseek-ai/DeepEP#65

No description provided.

Rows
Columns