deepseek-ai/DeepEP
1
0
Fork 0
mirror of https://github.com/deepseek-ai/DeepEP.git synced 2026-11-05 00:47:32 +00:00
Code Issues 308 Projects Packages Wiki Activity

[GH-ISSUE #427] low latency dispatch may have correctness issue under heavily imbalanced expert load cases #286

New issue
Closed
opened 2026-10-26 09:58:00 +00:00 by bsakaguchi · 11 comments
bsakaguchi commented 2026-10-26 09:58:00 +00:00
Owner
Copy link

Originally created by @yifeizhang-c on GitHub (Sep 24, 2025).
Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/427

The issue happens on the latest main da6ca24e as well as previous commit 8b0c5944, seems to be a long existing issue.

If we test test_low_latency.py on multiple nodes (e.g. 2 nodes, 16 ranks in total) with a larger num_tokens value (e.g. 2304) while making expert load heavily imbalanced (e.g. by changing the random scores into the following

    scores = torch.randn((num_tokens, num_experts), dtype=torch.float32, device='cuda').abs() + 1
    scores_imbalanced = torch.randn((1, num_experts), dtype=torch.float32, device='cuda').abs() + 1
    scores[10:, :] = scores_imbalanced.repeat(num_tokens - 10, 1)

):

The test might fail at dispatch stage correctness check, e.g.

  File "xxxxx", line 84, in test_main
    assert num_valid_tokens == (all_topk_idx == expert_id).sum().item(), f'{num_valid_tokens} != {(all_topk_idx == expert_id).sum().item()}'
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
AssertionError: 9181 != 6888

A more detailed analysis showing in above mentioned cases some inter-node ibgda num_tokens transmissions might have value mismatch, e.g.

[RANK 9] Performing nvshmemi_ibgda_amo_nonfetch_add, sm_id: 26, dst_expert_local_idx: 1, dst_rank: 5, num_tokens_sent: 2292
...
[RANK 5] sm_id: 46, local_expert_idx: 1, src_rank: 9, num_recv_tokens: 4585

This kind of failure often appears as the transmitted -num_tokens_sent - 1 got doubled or tripled in its value on arrival.

The above shown case is one of the low_latency dispatch correctness issues that we may encounter. Similar ibgda related transmissions issues may also lead to recv_x content different from expectation.

  File "xxx", line 101, in test_main
    assert (recv_x_amin == j - rank_offset).sum().item() == (all_topk_idx[j] == expert_id).sum().item()
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
AssertionError
Originally created by @yifeizhang-c on GitHub (Sep 24, 2025). Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/427 The issue happens on the latest main [da6ca24e](https://github.com/deepseek-ai/DeepEP/tree/da6ca24ef7c2594e058def142adfb8f2e31a7f85) as well as previous commit [8b0c5944](https://github.com/deepseek-ai/DeepEP/tree/8b0c59448bd808366d16fbcf036a231de9c5d0d2), seems to be a long existing issue. If we test `test_low_latency.py` on **multiple nodes** (e.g. 2 nodes, 16 ranks in total) with **a larger `num_tokens` value** (e.g. 2304) while making expert load **heavily imbalanced** (e.g. by changing the random scores into the following ``` scores = torch.randn((num_tokens, num_experts), dtype=torch.float32, device='cuda').abs() + 1 scores_imbalanced = torch.randn((1, num_experts), dtype=torch.float32, device='cuda').abs() + 1 scores[10:, :] = scores_imbalanced.repeat(num_tokens - 10, 1) ``` ): The test might fail at dispatch stage correctness check, e.g. ``` File "xxxxx", line 84, in test_main assert num_valid_tokens == (all_topk_idx == expert_id).sum().item(), f'{num_valid_tokens} != {(all_topk_idx == expert_id).sum().item()}' ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ AssertionError: 9181 != 6888 ``` A more detailed analysis showing in above mentioned cases some inter-node ibgda num_tokens transmissions might have value mismatch, e.g. ``` [RANK 9] Performing nvshmemi_ibgda_amo_nonfetch_add, sm_id: 26, dst_expert_local_idx: 1, dst_rank: 5, num_tokens_sent: 2292 ... [RANK 5] sm_id: 46, local_expert_idx: 1, src_rank: 9, num_recv_tokens: 4585 ``` This kind of failure often appears as the transmitted `-num_tokens_sent - 1` got doubled or tripled in its value on arrival. The above shown case is one of the low_latency dispatch correctness issues that we may encounter. Similar ibgda related transmissions issues may also lead to `recv_x` content different from expectation. ``` File "xxx", line 101, in test_main assert (recv_x_amin == j - rank_offset).sum().item() == (all_topk_idx[j] == expert_id).sum().item() ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ AssertionError ```
bsakaguchi commented 2026-10-26 09:58:01 +00:00
Author
Owner
Copy link

@sphish commented on GitHub (Sep 25, 2025):

I discussed this issue with @yifeizhang-c on WeChat, and we found that the cause is that in cases with a large number of tokens and a highly uneven distribution, a higher NVSHMEM_IB_QP_DEPTH needs to be set. I will consider how to make this problem easier to debug.

<!-- gh-comment-id:3332176678 --> @sphish commented on GitHub (Sep 25, 2025): I discussed this issue with @yifeizhang-c on WeChat, and we found that the cause is that in cases with a large number of tokens and a highly uneven distribution, a higher `NVSHMEM_IB_QP_DEPTH` needs to be set. I will consider how to make this problem easier to debug.
bsakaguchi commented 2026-10-26 09:58:01 +00:00
Author
Owner
Copy link

@sphish commented on GitHub (Sep 25, 2025):

Added a related assert. https://github.com/deepseek-ai/DeepEP/pull/430

<!-- gh-comment-id:3333057311 --> @sphish commented on GitHub (Sep 25, 2025): Added a related assert. https://github.com/deepseek-ai/DeepEP/pull/430
bsakaguchi commented 2026-10-26 09:58:02 +00:00
Author
Owner
Copy link

@yifeizhang-c commented on GitHub (Sep 25, 2025):

Thanks for the clarification! Enhancing the value of NVSHMEM_QP_DEPTH can effectively resolve all encountered low latency hanging/illegal memory access/incorrectness cases so far. Also thanks for the new enhancement for checking NVSHMEM_QP_DEPTH.

<!-- gh-comment-id:3333147018 --> @yifeizhang-c commented on GitHub (Sep 25, 2025): Thanks for the clarification! Enhancing the value of `NVSHMEM_QP_DEPTH` can effectively resolve all encountered low latency hanging/illegal memory access/incorrectness cases so far. Also thanks for the new enhancement for checking `NVSHMEM_QP_DEPTH`.
bsakaguchi commented 2026-10-26 09:58:02 +00:00
Author
Owner
Copy link

@yifeizhang-c commented on GitHub (Sep 26, 2025):

Hi @sphish , I would like to double confirm for the safe lower bound of NVSHMEM_QP_DEPTH, as large QP_DEPTH may lead to insufficient memory and cause qp creation failure.

Currently, the lower bound is set to be (num_max_dispatch_tokens_per_rank + 1) * 2. Whether it would be better to set it as num_max_dispatch_tokens_per_rank + num_ranks as the number of token-sending requests to each experts is at most num_max_dispatch_tokens_per_rank, and the number of counts/flag related sync request is at most num_ranks per experts?

<!-- gh-comment-id:3336494486 --> @yifeizhang-c commented on GitHub (Sep 26, 2025): Hi @sphish , I would like to double confirm for the safe lower bound of `NVSHMEM_QP_DEPTH`, as large `QP_DEPTH` may lead to insufficient memory and cause qp creation failure. Currently, the lower bound is set to be `(num_max_dispatch_tokens_per_rank + 1) * 2`. Whether it would be better to set it as `num_max_dispatch_tokens_per_rank + num_ranks` as the number of token-sending requests to each experts is at most `num_max_dispatch_tokens_per_rank`, and the number of counts/flag related sync request is at most `num_ranks` per experts?
bsakaguchi commented 2026-10-26 09:58:02 +00:00
Author
Owner
Copy link

@sphish commented on GitHub (Sep 26, 2025):

@yifeizhang-c Each QP can send at most num_max_dispatch_tokens_per_rank + 1 requests per round (since each expert has its own QP, each QP only sends a single flag).

Regarding the *2, it’s somewhat similar to a double-buffer design — this is because when we enter the current communication round, we cannot guarantee that all requests from the previous round have already been fully sent.

For example, using the test case you provided, if we set num_tokens = 864 and qp_depth = 1024, there is still a high probability of reproducing the issue.

<!-- gh-comment-id:3336681692 --> @sphish commented on GitHub (Sep 26, 2025): @yifeizhang-c Each QP can send at most `num_max_dispatch_tokens_per_rank + 1` requests per round (since each expert has its own QP, each QP only sends a single flag). Regarding the `*2`, it’s somewhat similar to a double-buffer design — this is because when we enter the current communication round, we cannot guarantee that all requests from the previous round have already been fully sent. For example, using the test case you provided, if we set `num_tokens = 864` and `qp_depth = 1024`, there is still a high probability of reproducing the issue.
bsakaguchi commented 2026-10-26 09:58:03 +00:00
Author
Owner
Copy link

@yifeizhang-c commented on GitHub (Sep 26, 2025):

Thanks for the clarification!

<!-- gh-comment-id:3336693224 --> @yifeizhang-c commented on GitHub (Sep 26, 2025): Thanks for the clarification!
bsakaguchi commented 2026-10-26 09:58:03 +00:00
Author
Owner
Copy link

@robertgshaw2-redhat commented on GitHub (Sep 26, 2025):

Would this issue manifest as a hang in the dispatch?

Thread 1 "VLLM::EngineCor" stopped.
[Switching focus to CUDA kernel 0, grid 1127099, cluster (32,0,0), block (64,0,0), thread (0,0,0), device 0, sm 58, warp 0, lane 0]
0x0000792b51cb3450 in void deep_ep::internode_ll::dispatch<true, false, 7168>(void*, void*, int*, long*, int*, int*, long*, void*, int*, void*, void const*, long const*, int*, int*, int*, int, int, int, int, int, int, int, int, int, bool, int)<<<(2,1,1),(1024,1,1)>>> ()

I'm getting a hang in multi-node case on B200

I upped the NVSHMEM_QP_DEPTH to 8192

note: this issue occurs on the most recent tagged release

<!-- gh-comment-id:3338912995 --> @robertgshaw2-redhat commented on GitHub (Sep 26, 2025): Would this issue manifest as a hang in the dispatch? ``` Thread 1 "VLLM::EngineCor" stopped. [Switching focus to CUDA kernel 0, grid 1127099, cluster (32,0,0), block (64,0,0), thread (0,0,0), device 0, sm 58, warp 0, lane 0] 0x0000792b51cb3450 in void deep_ep::internode_ll::dispatch<true, false, 7168>(void*, void*, int*, long*, int*, int*, long*, void*, int*, void*, void const*, long const*, int*, int*, int*, int, int, int, int, int, int, int, int, int, bool, int)<<<(2,1,1),(1024,1,1)>>> () ``` I'm getting a hang in multi-node case on B200 I upped the `NVSHMEM_QP_DEPTH` to 8192 note: this issue occurs on the most recent tagged release
bsakaguchi commented 2026-10-26 09:58:04 +00:00
Author
Owner
Copy link

@robertgshaw2-redhat commented on GitHub (Sep 26, 2025):

Would this issue manifest as a hang in the dispatch?

Thread 1 "VLLM::EngineCor" stopped.
[Switching focus to CUDA kernel 0, grid 1127099, cluster (32,0,0), block (64,0,0), thread (0,0,0), device 0, sm 58, warp 0, lane 0]
0x0000792b51cb3450 in void deep_ep::internode_ll::dispatch<true, false, 7168>(void*, void*, int*, long*, int*, int*, long*, void*, int*, void*, void const*, long const*, int*, int*, int*, int, int, int, int, int, int, int, int, int, bool, int)<<<(2,1,1),(1024,1,1)>>> ()

I'm getting a hang in multi-node case on B200

I upped the NVSHMEM_QP_DEPTH to 8192

note: this issue occurs on the most recent tagged release

apologies, I realized the other node is actually hung in cutlass MLA --- this is not due to the deepep.

[Switching focus to CUDA kernel 0, grid 1128202, cluster (37,0,0), block (74,0,0), thread (0,0,0), device 0, sm 68, warp 0, lane 0]
0x00007fb777f21120 in void cutlass::device_kernel<cutlass::fmha::kernel::Sm100FmhaMlaKernelTmaWarpspecialized<cute::tuple<cute::C<128>, cute::C<128>, cute::tuple<cute::C<512>, cute::C<64> > >, cutlass::bfloat16_t, float, cutlass::bfloat16_t, float, cutlass::fmha::kernel::Sm100MlaPersistentTileScheduler, false> >(cutlass::fmha::kernel::Sm100FmhaMlaKernelTmaWarpspecialized<cute::tuple<cute::C<128>, cute::C<128>, cute::tuple<cute::C<512>, cute::C<64> > >, cutlass::bfloat16_t, float, cutlass::bfloat16_t, float, cutlass::fmha::kernel::Sm100MlaPersistentTileScheduler, false>::Params)<<<(148,1,1),(256,1,1)>>> ()
<!-- gh-comment-id:3338922362 --> @robertgshaw2-redhat commented on GitHub (Sep 26, 2025): > Would this issue manifest as a hang in the dispatch? > > ``` > Thread 1 "VLLM::EngineCor" stopped. > [Switching focus to CUDA kernel 0, grid 1127099, cluster (32,0,0), block (64,0,0), thread (0,0,0), device 0, sm 58, warp 0, lane 0] > 0x0000792b51cb3450 in void deep_ep::internode_ll::dispatch<true, false, 7168>(void*, void*, int*, long*, int*, int*, long*, void*, int*, void*, void const*, long const*, int*, int*, int*, int, int, int, int, int, int, int, int, int, bool, int)<<<(2,1,1),(1024,1,1)>>> () > ``` > > I'm getting a hang in multi-node case on B200 > > I upped the `NVSHMEM_QP_DEPTH` to 8192 > > note: this issue occurs on the most recent tagged release apologies, I realized the other node is actually hung in cutlass MLA --- this is not due to the deepep. ```bash [Switching focus to CUDA kernel 0, grid 1128202, cluster (37,0,0), block (74,0,0), thread (0,0,0), device 0, sm 68, warp 0, lane 0] 0x00007fb777f21120 in void cutlass::device_kernel<cutlass::fmha::kernel::Sm100FmhaMlaKernelTmaWarpspecialized<cute::tuple<cute::C<128>, cute::C<128>, cute::tuple<cute::C<512>, cute::C<64> > >, cutlass::bfloat16_t, float, cutlass::bfloat16_t, float, cutlass::fmha::kernel::Sm100MlaPersistentTileScheduler, false> >(cutlass::fmha::kernel::Sm100FmhaMlaKernelTmaWarpspecialized<cute::tuple<cute::C<128>, cute::C<128>, cute::tuple<cute::C<512>, cute::C<64> > >, cutlass::bfloat16_t, float, cutlass::bfloat16_t, float, cutlass::fmha::kernel::Sm100MlaPersistentTileScheduler, false>::Params)<<<(148,1,1),(256,1,1)>>> () ```
bsakaguchi commented 2026-10-26 09:58:04 +00:00
Author
Owner
Copy link

@aahouzi commented on GitHub (Oct 3, 2025):

@yifeizhang-c @sphish I noticed that even after your fix for NVSHMEM_QP_DEPTH, I still have corectness issues as I noticed it never works for number of tokens larger than 512:

Traceback (most recent call last):
  File "/home/dpsk_a2a/DeepEP/tests/test_low_latency.py", line 257, in <module>
    torch.multiprocessing.spawn(test_loop, args=(num_processes, args), nprocs=num_processes)
  File "/usr/local/lib/python3.12/dist-packages/torch/multiprocessing/spawn.py", line 340, in spawn
    return start_processes(fn, args, nprocs, join, daemon, start_method="spawn")
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/usr/local/lib/python3.12/dist-packages/torch/multiprocessing/spawn.py", line 296, in start_processes
    while not context.join():
              ^^^^^^^^^^^^^^
  File "/usr/local/lib/python3.12/dist-packages/torch/multiprocessing/spawn.py", line 215, in join
    raise ProcessRaisedException(msg, error_index, failed_process.pid)
torch.multiprocessing.spawn.ProcessRaisedException: 

-- Process 6 terminated with the following error:
Traceback (most recent call last):
  File "/usr/local/lib/python3.12/dist-packages/torch/multiprocessing/spawn.py", line 90, in _wrap
    fn(i, *args)
  File "/home/dpsk_a2a/DeepEP/tests/test_low_latency.py", line 213, in test_loop
    test_main(num_tokens, hidden, num_experts, num_topk, rank, num_ranks, group, buffer,
  File "/home/dpsk_a2a/DeepEP/tests/test_low_latency.py", line 97, in test_main
    assert calc_diff(recv_x[:, -1], recv_src_info.view(-1)) < 0.007
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
AssertionError
<!-- gh-comment-id:3366225227 --> @aahouzi commented on GitHub (Oct 3, 2025): @yifeizhang-c @sphish I noticed that even after your fix for NVSHMEM_QP_DEPTH, I still have corectness issues as I noticed it never works for number of tokens larger than 512: ``` Traceback (most recent call last): File "/home/dpsk_a2a/DeepEP/tests/test_low_latency.py", line 257, in <module> torch.multiprocessing.spawn(test_loop, args=(num_processes, args), nprocs=num_processes) File "/usr/local/lib/python3.12/dist-packages/torch/multiprocessing/spawn.py", line 340, in spawn return start_processes(fn, args, nprocs, join, daemon, start_method="spawn") ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/usr/local/lib/python3.12/dist-packages/torch/multiprocessing/spawn.py", line 296, in start_processes while not context.join(): ^^^^^^^^^^^^^^ File "/usr/local/lib/python3.12/dist-packages/torch/multiprocessing/spawn.py", line 215, in join raise ProcessRaisedException(msg, error_index, failed_process.pid) torch.multiprocessing.spawn.ProcessRaisedException: -- Process 6 terminated with the following error: Traceback (most recent call last): File "/usr/local/lib/python3.12/dist-packages/torch/multiprocessing/spawn.py", line 90, in _wrap fn(i, *args) File "/home/dpsk_a2a/DeepEP/tests/test_low_latency.py", line 213, in test_loop test_main(num_tokens, hidden, num_experts, num_topk, rank, num_ranks, group, buffer, File "/home/dpsk_a2a/DeepEP/tests/test_low_latency.py", line 97, in test_main assert calc_diff(recv_x[:, -1], recv_src_info.view(-1)) < 0.007 ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ AssertionError ```
bsakaguchi commented 2026-10-26 09:58:04 +00:00
Author
Owner
Copy link

@sphish commented on GitHub (Oct 10, 2025):

@aahouzi This is caused by round_scale. round_scale will introduce relatively large error. In our test script, we set an upper bound for our commonly used test cased, but this may not be suitable for large num_tokens. If you don't use round_scale, you can disable the round_scale test.

<!-- gh-comment-id:3388315736 --> @sphish commented on GitHub (Oct 10, 2025): @aahouzi This is caused by `round_scale`. `round_scale` will introduce relatively large error. In our test script, we set an upper bound for our commonly used test cased, but this may not be suitable for large `num_tokens`. If you don't use `round_scale`, you can disable the `round_scale` test.
bsakaguchi commented 2026-10-26 09:58:04 +00:00
Author
Owner
Copy link

@gangxie112 commented on GitHub (Oct 11, 2025):

@sphish Hi, could this issue exist in normal kernel?

<!-- gh-comment-id:3393239732 --> @gangxie112 commented on GitHub (Oct 11, 2025): @sphish Hi, could this issue exist in normal kernel?
Sign in to join this conversation.
No Branch/Tag specified
Branches Tags
main
hybrid-ep
chore/remove-clang-format-yapf
v2.5
nvDev
try-fix-cu128
mori-ep
antgroup-opt
compat-old-torch
add-linter
eager-rdma
v1.2.1
Labels
Clear labels   
pull-request
Mirrored from GitHub Pull Request

No labels
pull-request
Milestone
Clear milestone
No items
No milestone
Projects
Clear projects
No items
No project
Assignees
Clear assignees
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#286
No description provided.