deepseek-ai/DeepGEMM
1
0
Fork 0
mirror of https://github.com/deepseek-ai/DeepGEMM.git synced 2026-11-05 00:47:26 +00:00
Code Issues 162 Projects Packages Wiki Activity

[GH-ISSUE #52] [Question] Why use one Math Warp-Group when block_m==64? #31

New issue
Closed
opened 2026-10-26 21:10:55 +00:00 by bsakaguchi · 1 comment
bsakaguchi commented 2026-10-26 21:10:55 +00:00
Owner
Copy link

Originally created by @ZJUGuoShuai on GitHub (Mar 10, 2025).
Original GitHub issue: https://github.com/deepseek-ai/DeepGEMM/issues/52

First of all, thanks to all of you guys open-sourcing this amazing repo for showing us how to implement an optimized FP8 GEMM.

I've been reading the implementation these days for learning. And I wonder why it launches CTAs with 1 Math Warp-Group when block_m==64 while launching CTAs with 2 Math Warp-Groups otherwise?

This is the code that determines how many Math Warp-Groups a CTA has, based on block_m.

template <uint32_t kNumTMAThreads, uint32_t kNumMathThreadsPerGroup>
__device__ __host__ constexpr int get_num_threads_per_sm(int block_m) {
    DG_STATIC_ASSERT(kNumMathThreadsPerGroup == 128, "Only support 128 threads per math group");
    return (block_m == 64 ? 1 : 2) * kNumMathThreadsPerGroup + kNumTMAThreads;
}

Can someone give me some hints?

Originally created by @ZJUGuoShuai on GitHub (Mar 10, 2025). Original GitHub issue: https://github.com/deepseek-ai/DeepGEMM/issues/52 First of all, thanks to all of you guys open-sourcing this amazing repo for showing us how to implement an optimized FP8 GEMM. I've been reading the implementation these days for learning. And I wonder **why it launches CTAs with 1 Math Warp-Group when `block_m==64` while launching CTAs with 2 Math Warp-Groups otherwise**? This is the code that determines how many Math Warp-Groups a CTA has, based on `block_m`. ```cpp template <uint32_t kNumTMAThreads, uint32_t kNumMathThreadsPerGroup> __device__ __host__ constexpr int get_num_threads_per_sm(int block_m) { DG_STATIC_ASSERT(kNumMathThreadsPerGroup == 128, "Only support 128 threads per math group"); return (block_m == 64 ? 1 : 2) * kNumMathThreadsPerGroup + kNumTMAThreads; } ``` Can someone give me some hints?
bsakaguchi commented 2026-10-26 21:10:57 +00:00
Author
Owner
Copy link

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

The m dim of WGMMA instruction is 64 (no m-32 instructions), see the PTX doc: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-shape

<!-- gh-comment-id:2709454331 --> @LyricZhao commented on GitHub (Mar 10, 2025): The m dim of WGMMA instruction is 64 (no m-32 instructions), see the PTX doc: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-shape
Sign in to join this conversation.
No Branch/Tag specified
Branches Tags
main
public-sync-20261014
0930-release
nv_dev
fix-mega-moe-fence
public-release-2609
public-release-2607
public-release-2606
mega-update
public-release-260416
ko3n1g/ci/build-wheels
nv_dev_f8e8fb5
nv_dev_950d31a
nv_dev_67fc648
nv_dev_c491439
nv_dev_4ff3f54
v2.1.1.post3
v2.1.1.post2
v2.1.1.post1
v2.1.1
v2.1.0
Labels
Clear labels   
pull-request
Mirrored from GitHub Pull Request

No labels
pull-request
Milestone
Clear milestone
No items
No milestone
Projects
Clear projects
No items
No project
Assignees
Clear assignees
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/DeepGEMM#31
No description provided.