deepseek-ai / deepseek-ai/DeepEP
Question about the global synchronization/barrier in internode notify_dispatch
- Dominant language
- Cuda
- Stars
- 10.1k
- Forks
- 1.4k
- Avg merge
- 4d 1h
- Merged PRs (30d)
- 2
Description
Hi, deepep team.
#### There are 3 global synchronization/barrier in internode notify_dispatch.
1. At the beginning of the notify_dispatch, SM0 performs a global barrier (the first warp do intra-node sync, the second warp do internode sync)
2. After calling nvshmem_int_put_nbi() to send numbers of tokens per rank/expert to RDMA ranks, SM0 performs a global barrier again.
3. At the end of the notify_dispatch, SM0 performs the last global barrier.
I roughly understand that the purpose of the **second** global barrier is to ensure all the rdma operations have been successfully completed and the data has been written to the memory of the destination rdma ranks.
So what is the purpose of the other two global barrier? one is at the beginning of the function, and the other is at the end.
As far as I know, a global barrier in nvshmem is a collective communication, and is time-consuming. In the profile file you provided, it also proves that notify_dispatch takes about 1ms.
So my question is, is there a chance to remove the global barrier at the beginning and end of the notify_dispatch?
@LyricZhao @sphish friendly ping
Contributor guide
No contributing guide indexed for this repository
Assessment
This issue has not been assessed yet.