NVIDIA / NVIDIA/cutlass

[QST] How SFB Is Organized When M=128 with 2-CTA Mode?

Open
#3,046 2 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

? - Needs Triage inactive-30d inactive-90d question
Dominant language
C++
Stars
10.5k
Forks
2.1k
Avg merge
3d 11h
Merged PRs (30d)
7

Description

The PTX doc is not very clear about SFB data organization on tensor memory when M=128 and 2-CTA mode. In CUTLASS, the S2T copy instruction for SFA/B always use tcgen05.cp.32x128b.warp4 no matter what the TmemAllocMode is.

I have two questions: how the ScaleFactorDuplicated2by2 affects the S2T process given the same 32x128b instruction used? What's the purpose of tcgen05.cp.64x128b.warp2::02_13 since CUTLASS only wraps it and not uses it?

@ccecka @thakkarV

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 with the PTX documentation on SFB organization for M=128 in 2-CTA mode, then compare ScaleFactorDuplicated2by2 with CUTLASS's tcgen05.cp.32x128b.warp4 S2T usage. Investigate the documented purpose of tcgen05.cp.64x128b.warp2::02_13 and whether CUTLASS's wrapper provides relevant context. Done means both questions have an authoritative explanation, with any needed documentation clarification identified.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
documentation, hpc
Issue type
Documentation
Difficulty
4/5
Estimated time
3-5 days
Activity status
Quiet
Clarity
Mostly clear
Newbie friendliness
42/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.