deepseek-ai / deepseek-ai/DeepEP

Question about the global synchronization/barrier in internode notify_dispatch

Open
#94 7 comments 0 reactions 0 assignees View on GitHub
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.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.