deepseek-ai/DeepEP

Fork 0

mirror of https://github.com/deepseek-ai/DeepEP.git synced 2026-11-05 00:47:32 +00:00

[GH-ISSUE #577] Question about synchronization guarantee in DeepEP dispatch low latency kernel #355

New issue

Open

opened 2026-10-26 09:59:02 +00:00 by bsakaguchi · 3 comments

bsakaguchi commented

2026-10-26 09:59:02 +00:00

Owner

Originally created by @Yi-sir on GitHub (Feb 2, 2026).
Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/577

Hi,
I've been reading the DeepEP dispatch low latency kernel code, but I have a concern about the synchronization mechanism:
My understanding of the issue:
In the dispatch kernel, each rank can only guarantee that the tokens it sends have been received by the target expert (through the atomic_finish_counter_per_expert mechanism). However, I don't see a synchronization mechanism that ensures the kernel exits only after receiving all expected tokens from other ranks.

Specifically:
The sending side uses atomic_finish_counter_per_expert to track local send completion
The receiving side uses while (ld_acquire_sys_global(rdma_recv_count...) == 0) to wait for notifications from specific ranks
But there seems to be no barrier or counter that aggregates "total expected receives from all ranks" to ensure the kernel waits for all of them before returning

My question:
Is there a mechanism I'm missing that guarantees a rank has received tokens from all other ranks before the kernel exits? Or does the kernel actually exit without ensuring full reception, relying on external synchronization (e.g., nvshmem_barrier_all() or a second kernel launch) to guarantee consistency?

I would appreciate any clarification on this. Thank you very much!

Originally created by @Yi-sir on GitHub (Feb 2, 2026). Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/577 Hi, I've been reading the DeepEP dispatch low latency kernel code, but I have a concern about the synchronization mechanism: My understanding of the issue: In the dispatch kernel, each rank can only guarantee that the tokens it sends have been received by the target expert (through the atomic_finish_counter_per_expert mechanism). However, I don't see a synchronization mechanism that ensures the kernel exits only after receiving all expected tokens from other ranks. Specifically: The sending side uses atomic_finish_counter_per_expert to track local send completion The receiving side uses while (ld_acquire_sys_global(rdma_recv_count...) == 0) to wait for notifications from specific ranks But there seems to be no barrier or counter that aggregates "total expected receives from all ranks" to ensure the kernel waits for all of them before returning My question: Is there a mechanism I'm missing that guarantees a rank has received tokens from all other ranks before the kernel exits? Or does the kernel actually exit without ensuring full reception, relying on external synchronization (e.g., nvshmem_barrier_all() or a second kernel launch) to guarantee consistency? I would appreciate any clarification on this. Thank you very much!

bsakaguchi commented

2026-10-26 09:59:05 +00:00

Author

Owner

@Yi-sir commented on GitHub (Feb 2, 2026):

As an additional question: When running DeepSeek-V3.1 model using SGLang, I observed in the profiling traces that the dispatch low latency recv operation takes longer than dispatch ll send. However, DeepEP/tests/test_low_latency.py shows the opposite results (where send takes longer than recv). Is this behavior expected/normal?

@Yi-sir commented on GitHub (Feb 2, 2026): As an additional question: When running DeepSeek-V3.1 model using SGLang, I observed in the profiling traces that the dispatch low latency recv operation takes longer than dispatch ll send. However, DeepEP/tests/test_low_latency.py shows the opposite results (where send takes longer than recv). Is this behavior expected/normal?

bsakaguchi commented

2026-10-26 09:59:06 +00:00

Author

Owner

@polarstormx commented on GitHub (Feb 3, 2026):

In the recv phase, each warp group is responsible for one src rank of a local expert. If a src rank's data hasn't arrived, the corresponding warp group will keep waiting in the while loop, preventing the kernel from exiting. Perhaps you thought that one while loop is waiting for multiple senders?

@polarstormx commented on GitHub (Feb 3, 2026): In the recv phase, each warp group is responsible for one src rank of a local expert. If a src rank's data hasn't arrived, the corresponding warp group will keep waiting in the while loop, preventing the kernel from exiting. Perhaps you thought that one while loop is waiting for multiple senders?

bsakaguchi commented

2026-10-26 09:59:06 +00:00

Author

Owner

@Yi-sir commented on GitHub (Feb 3, 2026):

In the recv phase, each warp group is responsible for one src rank of a local expert. If a src rank's data hasn't arrived, the corresponding warp group will keep waiting in the while loop, preventing the kernel from exiting. Perhaps you thought that one while loop is waiting for multiple senders?在接收阶段，每个线程束组负责处理本地专家一个源等级的数据。如果某个源等级的数据尚未到达，对应线程束组会在 while 循环中持续等待，导致内核无法退出。您是否误以为单个 while 循环在等待多个发送方？

Thank you! I did misunderstand what warp groups do in the recv phase...

@Yi-sir commented on GitHub (Feb 3, 2026): > In the recv phase, each warp group is responsible for one src rank of a local expert. If a src rank's data hasn't arrived, the corresponding warp group will keep waiting in the while loop, preventing the kernel from exiting. Perhaps you thought that one while loop is waiting for multiple senders?在接收阶段，每个线程束组负责处理本地专家一个源等级的数据。如果某个源等级的数据尚未到达，对应线程束组会在 while 循环中持续等待，导致内核无法退出。您是否误以为单个 while 循环在等待多个发送方？ Thank you! I did misunderstand what warp groups do in the recv phase...

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#355

No description provided.

Rows
Columns