mirror of
https://github.com/deepseek-ai/DeepGEMM.git
synced 2026-11-05 00:47:26 +00:00
[GH-ISSUE #51] Why is the initial value of barrier related to multicast? #33
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#33
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 @dajunHuang on GitHub (Mar 9, 2025).
Original GitHub issue: https://github.com/deepseek-ai/DeepGEMM/issues/51
Thank you for your amazing work！
I'm trying to understand your implemeatation and find the following mbarrier initialization code:
github.com/deepseek-ai/DeepGEMM@4d4f2342fe/deep_gemm/include/deep_gemm/fp8_gemm.cuh (L118)Can you help me explain why the initial value of empty barrier is
kNumTMAMulticast * kNumMathThreads / 32，instead of a more intuitive value likekNumMathThreads / 32? why a single arrive operation of empty barrier should be distrubuted tokNumTMAMulticastdifferent ctas like below? When I set the initial value tokNumMathThreads / 32and remove the arugment ofarrive(), program gets stuck forever.github.com/deepseek-ai/DeepGEMM@4d4f2342fe/deep_gemm/include/deep_gemm/fp8_gemm.cuh (L240)I'm not very familiar with the relationship between barrier and its multicast feature. I'd appreciate it If you can give me some ideas for understanding.
@soundOfDestiny commented on GitHub (Mar 9, 2025):
Multicast will rewrite shared memory in two CTAs so we cannot issue a multicast until neither of the corresponding shared memory part in the CTA is in use. That's why both CTAs need to issue an empty barrier.