Question about synchronization guarantee in DeepEP dispatch low latency kernel

[Yi-sir]

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!

[Yi-sir]

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?

[polarstormx]

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?

[Yi-sir]

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...