mirror of
https://github.com/deepseek-ai/DeepGEMM.git
synced 2026-11-05 00:47:26 +00:00
[GH-ISSUE #368] question about tmem barrier #177
Labels
No labels
pull-request
No milestone
No project
No assignees
1 participant
Notifications
Due date
No due date set.
Dependencies
No dependencies set.
Reference
deepseek-ai/DeepGEMM#177
Loading…
Add table
Add a link
Reference in a new issue
No description provided.
Delete branch "%!s()"
Deleting a branch is permanent. Although the deleted branch may continue to exist for a short time before it actually gets removed, it CANNOT be undone in most cases. Continue?
Originally created by @wukong1992 on GitHub (Jun 26, 2026).
Original GitHub issue: https://github.com/deepseek-ai/DeepGEMM/issues/368
I have a question about the comment here:
github.com/deepseek-ai/DeepGEMM@54e2261240/deep_gemm/include/deep_gemm/impls/sm100_fp8_fp4_mega_moe.cuh (L269)Since TMEM is consumed by both CTAs, the next line uses 2 * kNumEpilogueThreads. Shouldn't this barrier therefore use arrive at all CTAs instead?
Conversely, for tmem_full_barrier above, it seems that only the leader CTA calls arrive. Is my understanding correct?
@Rachmanino commented on GitHub (Jun 29, 2026):
When both CTA complete tcgen05_ld from TMEM, they arrive on leader CTA's tmem_empty_barrier to signal that their TMEM are both free, and is ready for next tcgen05mma.
When leader CTA's mma warp issues tcgen05mma, it issues tcgen05 commit on both CTA's tmem_empty_barrier, so that they can wait on their own mbarrier for tcgen05mma completion.