Fix dispatch timeout problem: concurrent submissions of qp should be forbidden.#181
Merged
LyricZhao merged 1 commit intodeepseek-ai:mainfrom Jun 13, 2025
Merged
Conversation
Collaborator
|
Regarding the example you mentioned, the following code prevents the situation where the ring doorbell is executed before the previous WQEs has been fully written: DeepEP/csrc/kernels/ibgda_device.cuh Lines 155 to 157 in 9fe9021 So I think the issue might not be in this part? |
Collaborator
|
Additionally, I noticed that you fixed a bug in the notify_dispatch function, thank you very much. Perhaps you could submit this change as a separate PR? |
|
Contributor
Author
YES, You are right~I miss it. But this patch indeed solved the dispatch timeout problem in my environment, I used a script to test it repeatedly and the problem no longer occurred, maybe it solved some potential problems that I didn't know about. Firstly, I suspected that the queue length was not enough and I increased the QP_DEPTH, but it didn't work. What's more, I added Do you have any idea about this problem or have any suggestion to solve this problem？Thank you very much. |
Contributor
Author
不好意思，没太明白您的意思，可以再展开说说嘛？ 我原来理解有误，以为出现 |
wqe的内容会被写入队列里，这个没有问题。但是目前看起来不同的sm会可能拿到都可用的pid，比如SM0拿到1，SM1拿到2，ring doorbell的时候会有SM1先执行，SM2后执行，会有error cqe。之前看到的一个解释是atomic的结果在不同SM上可能会不同 |
Contributor
|
I do once in a while see deepep dispatch timeout when handling SGLang large-scale system, thus wondering whether this can fix that... |
Collaborator
|
This issue is quite strange. I think we need to clarify the root cause of the bug before merging this PR. However, you've provided a good approach for reproducing the issue. I'll try to investigate it when I have some available time. |
Contributor
Author
Yeah, we should clarify the root cause of the bug. I find the following routine for the bug:
Guidance for the root cause: why kRDMAAndNVLForwarder reads stale The proving code and the log are as follows: diff --git a/csrc/kernels/internode.cu b/csrc/kernels/internode.cu
index 809e2a2..c352eb7 100644
--- a/csrc/kernels/internode.cu
+++ b/csrc/kernels/internode.cu
@@ -11,6 +11,8 @@ namespace internode {
extern nvshmem_team_t cpu_rdma_team;
+__device__ int dbg_iter = 0;
+
template<int kNumThreads, int kNumExpertsPerSM, int kNumRanksPerSM>
__global__ void __launch_bounds__(kNumThreads, 1)
get_dispatch_layout(const int64_t* topk_idx,
@@ -462,6 +464,8 @@ constexpr int get_num_topk_rdma_ranks(int num_rdma_ranks) {
return num_rdma_ranks < 8 ? num_rdma_ranks : 8;
}
+__device__ volatile bool cached_in_dst_nvl_rank[2][12][8][500]; // [rdma_ranks][channel_num][gpu_num_in_node][max_received_token_number]
+
template <bool kLowLatencyMode, int kNumRDMARanks, bool kCachedMode,
int kNumDispatchRDMASenderWarps, int kNumTopkRDMARanks = get_num_topk_rdma_ranks(kNumRDMARanks)>
__global__ void __launch_bounds__(((kNumDispatchRDMASenderWarps + 1 + NUM_MAX_NVL_PEERS) * 32), 1)
@@ -493,6 +497,10 @@ dispatch(int4* recv_x, float* recv_x_scales, int64_t* recv_topk_idx, float* recv
EP_DEVICE_ASSERT(ibgda_get_state()->num_rc_per_pe >= num_channels);
+ if (sm_id == 0 && thread_id == 0) {
+ dbg_iter++;
+ }
+
const auto role_meta = [=]() -> std::pair<WarpRole, int> {
if (is_forwarder) {
if (warp_id < NUM_MAX_NVL_PEERS) {
@@ -868,6 +876,22 @@ dispatch(int4* recv_x, float* recv_x_scales, int64_t* recv_topk_idx, float* recv
if (not kCachedMode)
send_nvl_head[i * NUM_MAX_NVL_PEERS] = cached_head;
}
+
+ if (lane_id == src_rdma_rank) {
+ if (dbg_iter == 1) {
+ if (i < 500) {
+ cached_in_dst_nvl_rank[src_rdma_rank][channel_id][dst_nvl_rank][i] = is_in_dst_nvl_rank;
+ } else {
+ printf("warning!!!!: i %d is >= 500\n", i);
+ }
+ } else {
+ if (cached_in_dst_nvl_rank[src_rdma_rank][channel_id][dst_nvl_rank][i] != is_in_dst_nvl_rank) {
+ printf("warning!!!!: dbg_iter %d src_rdma_rank %d channel_id %d dst_nvl_rank %d i %d is_in_dst_nvl_rank %d != cached_in_dst_nvl_rank %d \n",
+ dbg_iter, src_rdma_rank, channel_id, dst_nvl_rank, i, is_in_dst_nvl_rank, cached_in_dst_nvl_rank[src_rdma_rank][channel_id][dst_nvl_rank][i]);
+ }
+ }
+ }
+
if (not is_in_dst_nvl_rank)
continue;
@@ -1010,8 +1034,9 @@ dispatch(int4* recv_x, float* recv_x_scales, int64_t* recv_topk_idx, float* recv
// Timeout check
if (clock64() - start_time > NUM_TIMEOUT_CYCLES) {
- printf("DeepEP dispatch NVL receiver timeout, channel: %d, RDMA: %d, nvl: %d, src NVL: %d, head: %d, tail: %d\n",
- channel_id, rdma_rank, nvl_rank, src_nvl_rank, cached_channel_head_idx, cached_channel_tail_idx);
+ printf("rank[%d] DeepEP dispatch NVL receiver timeout, channel: %d, RDMA: %d, nvl: %d, src NVL: %d, head: %d, tail: [%p]->[%d], num_tokens_to_recv: [%d], start_offset: %d, end_offset: %d\n"
, rank,
+ channel_id, rdma_rank, nvl_rank, src_nvl_rank, cached_channel_head_idx, nvl_channel_tail.buffer(), cached_channel_tail_idx,
+ num_tokens_to_recv, start_offset, end_offset);
trap();
}
}
diff --git a/tests/test_internode.py b/tests/test_internode.py
index e9b3d57..9248201 100644
--- a/tests/test_internode.py
+++ b/tests/test_internode.py
@@ -172,8 +172,8 @@ def test_main(num_sms: int, local_rank: int, num_local_ranks: int, num_ranks: in
best_time, best_results = 1e10, None
rdma_send_bytes = (dispatch_bf16_rdma_send_bytes * fp8_factor) if isinstance(current_x, tuple) else dispatch_bf16_rdma_send_bytes
nvl_recv_bytes = (dispatch_bf16_nvl_recv_bytes * fp8_factor) if isinstance(current_x, tuple) else dispatch_bf16_nvl_recv_bytes
- for nvl_chunk_size in range(4, 33, 4):
- for rdma_chunk_size in range(4, 33, 4):
+ for nvl_chunk_size in [1, 4, 1, 4, 1, 4, 1, 4]:
+ for rdma_chunk_size in [1, 4, 1, 4, 1, 4, 1, 4]:
config = deep_ep.Config(num_sms, nvl_chunk_size, nvl_buffer_size, rdma_chunk_size, rdma_buffer_size)
tune_args = {'x': current_x, 'handle': handle, 'config': config}
t = bench(lambda: buffer.dispatch(**tune_args))[0]log from node 1: |
Contributor
|
这个修改在我的环境中试了下，是好使的，就不会遇到超时了， |
|
yes this pr also works in my H20 node with IB network. |
Collaborator
|
Hello, @alpha-baby @wzc-wuzhicheng @MiterV1 I was able to reproduce this issue in our environment. I found that it seems to be related to warp synchronization. The changes in this branch https://github.com/deepseek-ai/DeepEP/tree/fix-warp-sync have resolved the problem on my env. Could you please try this modification on your env to see if it works for you as well? |
Contributor
Hi sphish, Node A sends ib_write traffic to Node B, utilizing all 8 NICs. Running test_internode between Node A and Node B then consistently results in a timeout error, as shown below: I have tried applying all patches from this PR, but unfortunately, the problem has not been resolved. Hope this additional information is helpful. |
Collaborator
|
@polarstormx What issues are you experiencing with the ib_write_bw test? Also, could you open a new ISSUE and include your entire error log? I suspect that the problem you're encountering is not the same as the one reported by others in this issue. |
Collaborator
|
@wzc-wuzhicheng I found that using one QP per SM can slightly improve performance in some cases. Therefore, whether or not there are other ways to address this issue, I will merge this PR. However, I prefer to let the sender SM use the Can you modify this? |
Contributor
I get the same error log like #137 #158. And I can also reproduce these errors by using 4NIC H20. The performance of these triggering conditions is similar: test internode hang with small nvl_chunk_size and rdma_chunk_size, and there are no errors when using ibrc. So I think it's very likely for the same reason. The reason for using ibwrite is the discovery of hang when passing kv cache. If necessary, I will open another issue. Here's the Complete error log. The error in several scenarios is the same. |
Collaborator
|
@polarstormx I misunderstood what you said—I thought you were encountering issues when running |
Contributor
I tested this patch, and still reproduce timeout. node0 log: node1 log |
Collaborator
|
@alpha-baby Thank you for your feedback. It seems we are not encountering the same issue. |
|
@alpha-baby if you use a rocev2 underlay networking, pls checkout the switch config. |
Contributor
Author
I tested this patch, and still reproduce timeout. The |
Contributor
Author
Ok, is it necessary to submit a new PR to use one QP per SM and keep this PR for timeout discussion？ |
Contributor
What configuration does it specifically refer to? If there is something wrong with ROCE, that NVSHMEM should report the ver sbose error, but it doesn't, so I don't think there is anything wrong with the network. |
Collaborator
@wzc-wuzhicheng I think we can make the changes in the current PR. The timeout discussions can be moved to the issue thread. |
let the sender SM use the channel_id, and the receiver SM use channel_id + num_channels
1b6b611 to
6bd6585
Compare
Contributor
Author
How to move this pr to issue thread? |
Collaborator
|
@wzc-wuzhicheng I mean we can continue the discussion in this issue thread. |
6 tasks
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
8 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
kRDMASenderCoordinator and kForwarderCoordinator use the same qp to commit requests with channel_id, which will break down the atomicity of doorbell operation:
kRDMASenderCoordinator and kForwarderCoordinator will use different qp to commit requests with sm_id now.
NOTE: The timeout problem is not easy to reproduce in gpu:nic=1:1 environment , but I found that there are someways to reproduce it more easily:
related issues:
#137
#158