[BUG] SFA/SFB s2t (UTCCP) should be issued only by the leader CTA
Nobody has claimed this yet.
- Dominant language
- C++
- Stars
- 10.5k
- Forks
- 2.1k
- Avg merge
- 3d 11h
- Merged PRs (30d)
- 7
Description
Which component has the problem?
CUTLASS C++
Bug Report
Describe the bug
At cutlass/gemm/collective/sm100_blockscaled_mma_mixed_tma_cpasync_warpspecialized.hpp:L1029-1032, s2t copy is performed by both SMs when 2SM mode is enabled.
Steps/Code to reproduce bug
This itself seems to work correctly, but has a potential race condition bug.
Expected behavior
The stub should be protected by a branch that checks is_mma_lead_cta i.e.
if (is_mma_leader_cta && cute::elect_one_sync()) {
copy(tiled_copy_s2t_SFA, thr_tCsSFA_s2t(_,_,_,_,read_stage_tma), thr_tCtSFA_s2t);
copy(tiled_copy_s2t_SFB, thr_tCsSFB_s2t(_,_,_,_,read_stage_tma), thr_tCtSFB_s2t);
}
Environment details (please complete the following information):
- Environment location: B200
Contributor guide
No contributing guide indexed for this repository
First steps
- Read the whole issue, then the project's contributing guide.
- Comment on the issue to say you are picking it up — it saves two people doing the same work.
- Fork the repository and make your change on a branch.
- Open a pull request that references the issue number.
Research direction
Start at cutlass/gemm/collective/sm100_blockscaled_mma_mixed_tma_cpasync_warpspecialized.hpp:1029-1032 and trace the 2SM path around the s2t SFA/SFB copies. Check how is_mma_leader_cta and elect_one_sync are used nearby, then validate the B200 reproduction or relevant 2SM tests. Done means the copies are issued only by the leader CTA without introducing a race.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- performance
- Issue type
- Bug
- Difficulty
- 3/5
- Estimated time
- 1-2 days
- Activity status
- Quiet
- Clarity
- Clearly specified
- Newbie friendliness
- 68/100