[GH-ISSUE #44] Question about sm_id #40

New issue

Closed

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

bsakaguchi commented

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

Owner

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

In the source code, SETUP_LAUNCH_CONFIG(num_sms, kNumThreads, stream); sets num_sms as the parameter for gridDim, and uses auto sm_id = static_cast<int>(blockIdx.x); to map blocks to SMs one-to-one. I would like to ask how the code ensures that there is only one block per SM?

Originally created by @Qizhi697 on GitHub (Mar 5, 2025). Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/44 In the source code, `SETUP_LAUNCH_CONFIG(num_sms, kNumThreads, stream);` sets num_sms as the parameter for `gridDim`, and uses `auto sm_id = static_cast<int>(blockIdx.x);` to map blocks to SMs one-to-one. I would like to ask how the code ensures that there is only `one block per SM`?

bsakaguchi closed this issue

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

bsakaguchi commented

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

Author

Owner

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

This is a good question.

Actually, we didn't control this explicitly. But for most of the kernels in DeepEP, the registers are using too much, forcing a block must be in a SM, also see the __launch_bound__ attrs and PTXAS register usage output. Another explicit way to do this is allocating a large shared memory (but not use it), of which the size exceeds half of the physical shared memory limit.

Moreover, cudaLaunchAttributeCooperative ensures all blocks should be launched together, see SETUP_LAUNCH_CONFIG.

In terms of working with the computation kernels, you should make computation kernels persistent (e.g. compute 112 SMs, comm 132 - 112 = 20 SMs, total SMs is fixed, so the hardward scheduler will not make two blocks into one SM), this is also another way to ensure.

@LyricZhao commented on GitHub (Mar 5, 2025): This is a good question. Actually, we didn't control this explicitly. But for most of the kernels in DeepEP, the registers are using too much, forcing a block must be in a SM, also see the `__launch_bound__` attrs and PTXAS register usage output. Another explicit way to do this is allocating a large shared memory (but not use it), of which the size exceeds half of the physical shared memory limit. Moreover, `cudaLaunchAttributeCooperative` ensures all blocks should be launched together, see `SETUP_LAUNCH_CONFIG`. In terms of working with the computation kernels, you should make computation kernels persistent (e.g. compute 112 SMs, comm 132 - 112 = 20 SMs, total SMs is fixed, so the hardward scheduler will not make two blocks into one SM), this is also another way to ensure.

bsakaguchi commented

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

Author

Owner

@Qizhi697 commented on GitHub (Mar 5, 2025):

Thank you very much for your response. It was very clear, intuitive, and practical!

I have some questions about cudaLaunchAttributeCooperative. When launching all blocks, how does cudaLaunchAttributeCooperative differ from the regular cuda Launch Kernel method? What impact does cudaLaunchAttributeCooperative have on the allocation of blocks on SMs?

@Qizhi697 commented on GitHub (Mar 5, 2025): Thank you very much for your response. It was very clear, intuitive, and practical! I have some questions about `cudaLaunchAttributeCooperative`. When launching all blocks, how does `cudaLaunchAttributeCooperative` differ from the regular cuda Launch Kernel method? What impact does `cudaLaunchAttributeCooperative` have on the allocation of blocks on SMs?

bsakaguchi commented

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

Author

Owner

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

You can refer to https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__HIGHLEVEL.html#group__CUDART__HIGHLEVEL_1g7c4cb6c44a6c4608da36c44374499b31.

The biggest diff with the normal launch for me is that this function requires all blocks launch together.

For example, we want to launch 200 blocks, if one SM can hold 2 blocks, then it can be launched with cooperative API. If the reg/smem is not enough for 2 SMs, then the cooperative launch will raise an error (normal will not).

Only with cooperative launch, you can use the TMA multicast (a GEMM concept, multicasting some data to two blocks by reading once, you must ensure two blocks alive at the same time), or cg::this_grid().sync() (used in low-latency combine function). Imagine that if the hardware only has the resource to hold 132 blocks, but 200 in total, we can never do a whole grid sync (deadlock, the 200 - 132 blocks are not running and they cannot run because no regs/smem resources, that's why there's only __syncthreads but no __syncgrid before the cooperative launch CUDA API).

Or if another kernel on other stream is occuping 10 blocks (also SM), we want to launch 132 blocks (one SM cannot hold 2 blocks). I guess the hardware will wait until the 10 SMs are released.

@LyricZhao commented on GitHub (Mar 5, 2025): You can refer to https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__HIGHLEVEL.html#group__CUDART__HIGHLEVEL_1g7c4cb6c44a6c4608da36c44374499b31. The biggest diff with the normal launch for me is that this function requires **all blocks** launch together. For example, we want to launch 200 blocks, if one SM can hold 2 blocks, then it can be launched with cooperative API. If the reg/smem is not enough for 2 SMs, then the cooperative launch will raise an error (normal will not). Only with cooperative launch, you can use the TMA multicast (a GEMM concept, multicasting some data to two blocks by reading once, you must ensure two blocks alive at the same time), or `cg::this_grid().sync()` (used in low-latency combine function). Imagine that if the hardware only has the resource to hold 132 blocks, but 200 in total, we can never do a whole grid sync (deadlock, the 200 - 132 blocks are not running and they cannot run because no regs/smem resources, that's why there's only `__syncthreads` but no `__syncgrid` before the cooperative launch CUDA API). Or if another kernel on other stream is occuping 10 blocks (also SM), we want to launch 132 blocks (one SM cannot hold 2 blocks). I guess the hardware will wait until the 10 SMs are released.

bsakaguchi commented

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

Author

Owner

@Qizhi697 commented on GitHub (Mar 5, 2025):

Thank you for response, it was very clear! Thank you for your time!

@Qizhi697 commented on GitHub (Mar 5, 2025): Thank you for response, it was very clear! Thank you for your time!

bsakaguchi referenced this issue

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

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