[QST] wgmma.mma_async instructions are serialized due to insufficient register resources
Open
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
When I build the latest cutlass library for 90a, I see a lot of warnings like:
ptxas info : (C7511) Potential Performance Loss: wgmma.mma_async instructions are serialized due to insufficient register resources for the wgmma pipeline in the function '_ZN7cutlass13device_kernelI125cutlass3x_sm90_tensorop_s64x128x16gemm_bf16_bf16_f32_void_bf16_256x128x64_1x2x1_0_tnn_align8_warpspecialized_pingpong_epi_tmaEEvNT_6ParamsE'
Am I doing something wrong, or is that intended?
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 with the ptxas warning for the listed SM90 CUTLASS kernel function and reproduce the latest CUTLASS 90a build that emits it. Read the relevant WGMMA pipeline and register-resource documentation or implementation, then determine whether the serialization is expected; done means documenting the cause or identifying a concrete correction.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- performance
- Issue type
- Bug
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Stale
- Clarity
- Needs clarification
- Newbie friendliness
- 20/100