NVIDIA / NVIDIA/cutlass

[BUG] SFA/SFB s2t (UTCCP) should be issued only by the leader CTA

Open
#3,331 1 comment 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

? - Needs Triage bug CUTLASS C++ inactive-30d
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

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. 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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.