[GH-ISSUE #94] Question about the global synchronization/barrier in internode notify_dispatch

bsakaguchi commented

2026-10-26 09:54:00 +00:00

Owner

Originally created by @songhexiang on GitHub (Mar 26, 2025).
Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/94

Hi, deepep team.

There are 3 global synchronization/barrier in internode notify_dispatch.

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)
After calling nvshmem_int_put_nbi() to send numbers of tokens per rank/expert to RDMA ranks, SM0 performs a global barrier again.
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

Originally created by @songhexiang on GitHub (Mar 26, 2025). Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/94 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

bsakaguchi commented

2026-10-26 09:54:01 +00:00

Author

Owner

@LyricZhao commented on GitHub (Mar 26, 2025):

The first barrier ensures the last dispatch/combine is finished, as we are cleaning the head/tail/other metadata, otherwise the unfinished ranks will read illegal things.

The last barriers ensures the head/tail/other metadata are cleaned before actually running the dispatch/combine, otherwise the faster ranks will read uncleaned metadata.

There is some space to optimize for these parts, but may need some refactor, we plan to do it maybe several months later.

@LyricZhao commented on GitHub (Mar 26, 2025): The first barrier ensures the last dispatch/combine is finished, as we are cleaning the head/tail/other metadata, otherwise the unfinished ranks will read illegal things. The last barriers ensures the head/tail/other metadata are cleaned before actually running the dispatch/combine, otherwise the faster ranks will read uncleaned metadata. There is some space to optimize for these parts, but may need some refactor, we plan to do it maybe several months later.

bsakaguchi commented

2026-10-26 09:54:01 +00:00

Author

Owner

@songhexiang commented on GitHub (Mar 26, 2025):

Thanks for your prompt response!

@songhexiang commented on GitHub (Mar 26, 2025): Thanks for your prompt response!

bsakaguchi commented

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

Author

Owner

@songhexiang commented on GitHub (Mar 26, 2025):

Can I understand that notify_dispatch and dispatch share the same rdma buffer(void* rdma_buffer_ptr) and nvlink buffer(void** buffer_ptrs). These two global barriers are used to ensure that the rdma buffer and nvlink buffer are cleaned up for the next opertion.

@songhexiang commented on GitHub (Mar 26, 2025): Can I understand that notify_dispatch and dispatch share the same rdma buffer(void* rdma_buffer_ptr) and nvlink buffer(void** buffer_ptrs). These two global barriers are used to ensure that the rdma buffer and nvlink buffer are cleaned up for the next opertion.

bsakaguchi commented

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

Author

Owner

@LyricZhao commented on GitHub (Mar 26, 2025):

Yes, the data buffer is reused as you said. But another reason is that one of the notify's responsibility is to clean head/tail/metadata for later dispatch.

@LyricZhao commented on GitHub (Mar 26, 2025): Yes, the data buffer is reused as you said. But another reason is that one of the notify's responsibility is to clean head/tail/metadata for later dispatch.

bsakaguchi commented

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

Author

Owner

@Autumn1998 commented on GitHub (Apr 17, 2025):

Hi teams:

After the notify_dispatch, there is a CPU synchronize using wile loop which could cause problem when make DeepEP overlap with other compute kernels. what's more, sometimes notify_dispatch occupy several milliseconds which is much longer than normal (~100us), so I have some related questions:

the CPU sync is necessary here? any possibility to rm this sync?
why sometimes the notify_dispatch consumed more time than normal case?

Thanks~

@Autumn1998 commented on GitHub (Apr 17, 2025): Hi teams: After the `notify_dispatch`, there is a CPU synchronize using wile loop which could cause problem when make DeepEP overlap with other compute kernels. what's more, sometimes `notify_dispatch` occupy several milliseconds which is much longer than normal (~100us), so I have some related questions: - the CPU sync is necessary here? any possibility to rm this sync? - why sometimes the `notify_dispatch` consumed more time than normal case? Thanks~

bsakaguchi commented

2026-10-26 09:54:03 +00:00

Author

Owner

@terminator123 commented on GitHub (Oct 16, 2025):

Hi teams:

After the notify_dispatch, there is a CPU synchronize using wile loop which could cause problem when make DeepEP overlap with other compute kernels. what's more, sometimes notify_dispatch occupy several milliseconds which is much longer than normal (~100us), so I have some related questions:

the CPU sync is necessary here? any possibility to rm this sync?

why sometimes the notify_dispatch consumed more time than normal case?

Thanks~

same question，have you got the answer ?

@terminator123 commented on GitHub (Oct 16, 2025): > Hi teams: > > After the `notify_dispatch`, there is a CPU synchronize using wile loop which could cause problem when make DeepEP overlap with other compute kernels. what's more, sometimes `notify_dispatch` occupy several milliseconds which is much longer than normal (~100us), so I have some related questions: > > * the CPU sync is necessary here? any possibility to rm this sync? > * why sometimes the `notify_dispatch` consumed more time than normal case? > > Thanks~ same question，have you got the answer ?

bsakaguchi commented

2026-10-26 09:54:03 +00:00

Author

Owner

@xingzhu-HU commented on GitHub (Dec 17, 2025):

Hi teams: 各位团队成员：

After the notify_dispatch, there is a CPU synchronize using wile loop which could cause problem when make DeepEP overlap with other compute kernels. what's more, sometimes notify_dispatch occupy several milliseconds which is much longer than normal (~100us), so I have some related questions:在 notify_dispatch 之后，会有一个使用 while 循环的 CPU 同步，这可能会导致 DeepEP 与其他计算内核重叠时出现问题。此外， notify_dispatch 有时需要几毫秒的时间，远长于正常情况（约 100 微秒），因此我有一些相关问题：

the CPU sync is necessary here? any possibility to rm this sync?这里必须启用 CPU 同步吗？有没有办法取消这个同步？

why sometimes the notify_dispatch consumed more time than normal case?为什么有时 notify_dispatch 函数消耗的时间比正常情况要长？

Thanks~ 谢谢~

wait for answer too

@xingzhu-HU commented on GitHub (Dec 17, 2025): > Hi teams: 各位团队成员： > > After the `notify_dispatch`, there is a CPU synchronize using wile loop which could cause problem when make DeepEP overlap with other compute kernels. what's more, sometimes `notify_dispatch` occupy several milliseconds which is much longer than normal (~100us), so I have some related questions:在 `notify_dispatch` 之后，会有一个使用 while 循环的 CPU 同步，这可能会导致 DeepEP 与其他计算内核重叠时出现问题。此外， `notify_dispatch` 有时需要几毫秒的时间，远长于正常情况（约 100 微秒），因此我有一些相关问题： > > * the CPU sync is necessary here? any possibility to rm this sync?这里必须启用 CPU 同步吗？有没有办法取消这个同步？ > * why sometimes the `notify_dispatch` consumed more time than normal case?为什么有时 `notify_dispatch` 函数消耗的时间比正常情况要长？ > > Thanks~ 谢谢~ wait for answer too

bsakaguchi referenced this issue

2026-10-26 09:56:29 +00:00

[GH-ISSUE #270] A timeout error occurred when running test_internode.py on four devices. #200

Rows
Columns

[GH-ISSUE #94] Question about the global synchronization/barrier in internode notify_dispatch #86

There are 3 global synchronization/barrier in internode notify_dispatch.