[GH-ISSUE #92] Will the Normal Kernel (dispatch kernel) exclusively occupy the SM of the GPU? #84

Open
opened 2026-10-26 09:53:58 +00:00 by bsakaguchi · 1 comment
Owner

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

In the dispatch definition, each block has 512 threads. During actual execution, can we ensure that each block of the dispatch kernel exclusively occupies one SM through resource constraints such as register limits?

Originally created by @PhilipJohnson2 on GitHub (Mar 25, 2025). Original GitHub issue: https://github.com/deepseek-ai/DeepEP/issues/92 In the dispatch definition, each block has 512 threads. During actual execution, can we ensure that each block of the dispatch kernel exclusively occupies one SM through resource constraints such as register limits?
Author
Owner

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

Yes, the register usage is spilling for almost all kernels, and we also have things like __launch_bounds__(kNumThreads, 1) to tell 1 block per SM. So this is exclusive.

<!-- gh-comment-id:2753126055 --> @LyricZhao commented on GitHub (Mar 26, 2025): Yes, the register usage is spilling for almost all kernels, and we also have things like `__launch_bounds__(kNumThreads, 1)` to tell 1 block per SM. So this is exclusive.
Sign in to join this conversation.
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#84
No description provided.