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 #152] Question about overlapping between promotion and wgmma #88

New issue
Closed
opened 2026-10-26 21:12:13 +00:00 by bsakaguchi · 2 comments
bsakaguchi commented 2026-10-26 21:12:13 +00:00
Owner
Copy link

Originally created by @pengwubj on GitHub (Aug 4, 2025).
Original GitHub issue: https://github.com/deepseek-ai/DeepGEMM/issues/152

Why isn't a barrier used to establish a deterministic execution order between two warpgroups, so that promotion can be well hidden within the matrix computation? While studying CUTLASS, I found that the ping-pong approach uses OrderedSequenceBarrier to serialize the execution of two warpgroups, and for example, FA3 uses NamedBarrier to coordinate between two warpgroups（FLashMLA in same way）。Of course, it seems that this can also be automatically achieved based on NVIDIA hardware's greedy warp scheduling strategy and FFMA interleave technology.
Could you share your insight and experience about this ? Thanks a lot!

Originally created by @pengwubj on GitHub (Aug 4, 2025). Original GitHub issue: https://github.com/deepseek-ai/DeepGEMM/issues/152 Why isn't a barrier used to establish a deterministic execution order between two warpgroups, so that promotion can be well hidden within the matrix computation? While studying CUTLASS, I found that the ping-pong approach uses `OrderedSequenceBarrier` to serialize the execution of two warpgroups, and for example, FA3 uses `NamedBarrier` to coordinate between two warpgroups（FLashMLA in same way）。Of course, it seems that this can also be automatically achieved based on NVIDIA hardware's greedy warp scheduling strategy and FFMA interleave technology. Could you share your insight and experience about this ? Thanks a lot!
bsakaguchi commented 2026-10-26 21:12:14 +00:00
Author
Owner
Copy link

@LyricZhao commented on GitHub (Aug 14, 2025):

Barrier/sync costs a lot, we don't exchage data between two warpgroups. The overlapping is done by hardware warp scheduler.

<!-- gh-comment-id:3187571013 --> @LyricZhao commented on GitHub (Aug 14, 2025): Barrier/sync costs a lot, we don't exchage data between two warpgroups. The overlapping is done by hardware warp scheduler.
bsakaguchi commented 2026-10-26 21:12:15 +00:00
Author
Owner
Copy link

@pengwubj commented on GitHub (Aug 14, 2025):

Thanks a lot !

<!-- gh-comment-id:3187873880 --> @pengwubj commented on GitHub (Aug 14, 2025): Thanks a lot !
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#88
No description provided.