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 #61] Questions about Compute-Communication Overlap #58

New issue

Closed

opened 2026-10-26 09:53:32 +00:00 by bsakaguchi · 4 comments

bsakaguchi commented

2026-10-26 09:53:32 +00:00

Owner

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

hello, experts. see from codebase, I have several questions:

I think both async_finish and return_recv_hook can achieve compute-communication overlap, so what's the difference between them ?
with saync_finish mode, communication use comm_stream, it's make sense. but why return_recv_hook mode use compute_stream for communication ?

thanks.

Originally created by @Leaf996 on GitHub (Mar 9, 2025). Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/61 hello, experts. see from codebase, I have several questions: - I think both `async_finish` and `return_recv_hook` can achieve `compute-communication overlap`, so what's the difference between them ? - with `saync_finish` mode, communication use `comm_stream`, it's make sense. but why `return_recv_hook` mode use `compute_stream` for communication ? thanks.

bsakaguchi closed this issue

2026-10-26 09:53:32 +00:00

bsakaguchi commented

2026-10-26 09:53:33 +00:00

Author

Owner

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

They are two different ways to overlap:

async_finish: using two CUDA streams to overlap, computation and communication kernels share all the SMs (e.g. internode kernels using 20 SMs, and GEMM on the computation uses 132 - 20 = 112 SMs)
return_recv_hook: communication kernels does not use any SM, it just issue RDMA requests (as soon as possible, using more SMs, so just reuse the computation stream) and just return. The RDMA transmission is hid in the background.

For low-latency kernels, the async_finish is a legacy-style API which may not be used (to align with the inter/intra kernels), you can just ignore it. We have it, because when we didn't have the return_recv_hook API, we used it for shared experts overlapping (which is much easier, communication uses 96 SMs, shared GEMM uses 132 - 96 = 36 SMs), now we don't use it anymore.

@LyricZhao commented on GitHub (Mar 10, 2025): They are two different ways to overlap: - `async_finish`: using two CUDA streams to overlap, computation and communication kernels share all the SMs (e.g. internode kernels using 20 SMs, and GEMM on the computation uses 132 - 20 = 112 SMs) - `return_recv_hook`: communication kernels does not use any SM, it just issue RDMA requests (as soon as possible, using more SMs, so just reuse the computation stream) and just return. The RDMA transmission is hid in the background. For low-latency kernels, the `async_finish` is a legacy-style API which may not be used (to align with the inter/intra kernels), you can just ignore it. We have it, because when we didn't have the `return_recv_hook ` API, we used it for shared experts overlapping (which is much easier, communication uses 96 SMs, shared GEMM uses 132 - 96 = 36 SMs), now we don't use it anymore.

bsakaguchi commented

2026-10-26 09:53:33 +00:00

Author

Owner

@chenhongyu2048 commented on GitHub (Apr 21, 2025):

They are two different ways to overlap:

async_finish: using two CUDA streams to overlap, computation and communication kernels share all the SMs (e.g. internode kernels using 20 SMs, and GEMM on the computation uses 132 - 20 = 112 SMs)

return_recv_hook: communication kernels does not use any SM, it just issue RDMA requests (as soon as possible, using more SMs, so just reuse the computation stream) and just return. The RDMA transmission is hid in the background.

For low-latency kernels, the async_finish is a legacy-style API which may not be used (to align with the inter/intra kernels), you can just ignore it. We have it, because when we didn't have the return_recv_hook API, we used it for shared experts overlapping (which is much easier, communication uses 96 SMs, shared GEMM uses 132 - 96 = 36 SMs), now we don't use it anymore.

Hello, I would like to ask how to set the amount of SM used for gemm? Or do you think it's possible for me to find the answer in the DeepGEMM library? @LyricZhao Thanks for your reply!

@chenhongyu2048 commented on GitHub (Apr 21, 2025): > They are two different ways to overlap: > > * `async_finish`: using two CUDA streams to overlap, computation and communication kernels share all the SMs (e.g. internode kernels using 20 SMs, and GEMM on the computation uses 132 - 20 = 112 SMs) > * `return_recv_hook`: communication kernels does not use any SM, it just issue RDMA requests (as soon as possible, using more SMs, so just reuse the computation stream) and just return. The RDMA transmission is hid in the background. > > For low-latency kernels, the `async_finish` is a legacy-style API which may not be used (to align with the inter/intra kernels), you can just ignore it. We have it, because when we didn't have the `return_recv_hook ` API, we used it for shared experts overlapping (which is much easier, communication uses 96 SMs, shared GEMM uses 132 - 96 = 36 SMs), now we don't use it anymore. Hello, I would like to ask how to set the amount of SM used for gemm? Or do you think it's possible for me to find the answer in the DeepGEMM library? @LyricZhao Thanks for your reply!

bsakaguchi commented

2026-10-26 09:53:33 +00:00

Author

Owner

@LyricZhao commented on GitHub (Apr 22, 2025):

DeepGEMM: deep_gemm.set_num_sms(math_sms)
cuBLAS: CUBLAS_CHECK(cublasLtMatmulDescSetAttribute(desc, CUBLASLT_MATMUL_DESC_SM_COUNT_TARGET, &math_sms, sizeof(math_sms)));

@LyricZhao commented on GitHub (Apr 22, 2025): - DeepGEMM: `deep_gemm.set_num_sms(math_sms)` - cuBLAS: `CUBLAS_CHECK(cublasLtMatmulDescSetAttribute(desc, CUBLASLT_MATMUL_DESC_SM_COUNT_TARGET, &math_sms, sizeof(math_sms)));`

bsakaguchi commented

2026-10-26 09:53:34 +00:00

Author

Owner

@chenhongyu2048 commented on GitHub (Apr 22, 2025):

DeepGEMM: deep_gemm.set_num_sms(math_sms)

cuBLAS: CUBLAS_CHECK(cublasLtMatmulDescSetAttribute(desc, CUBLASLT_MATMUL_DESC_SM_COUNT_TARGET, &math_sms, sizeof(math_sms)));

Thanks for your help!

@chenhongyu2048 commented on GitHub (Apr 22, 2025): > * DeepGEMM: `deep_gemm.set_num_sms(math_sms)` > * cuBLAS: `CUBLAS_CHECK(cublasLtMatmulDescSetAttribute(desc, CUBLASLT_MATMUL_DESC_SM_COUNT_TARGET, &math_sms, sizeof(math_sms)));` Thanks for your help!

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

No description provided.

Rows
Columns