[QST] Why only N={64, 128, 192, 256} is supported for mxf8f6f4 GeMM?
Nobody has claimed this yet.
- Dominant language
- C++
- Stars
- 10.5k
- Forks
- 2.1k
- Avg merge
- 3d 11h
- Merged PRs (30d)
- 7
Description
What is your question?
I'm looking through the CUTLASS 3.9 source code and how gemm of mx types works on Blackwell hardware.
First off, I find tcgen05.mma.mxf8f6f4 supports {8,16,24,…,256} for 1SM mma and {16, 32, 48, ..., 256} for 2SM mma, which is reflected in CUTLASS API that those Ns are supported for non-blockscaled types (e.g., f8f6f4) in Collective builders. However, I can't see why Ns are restricted so only a few of those are supported like:
sm100_make_blockscaled_1sm_trivial_tiled_mma() {
// ...
// Do not allow a tiled MMA N mode > 1, as that is not reasonable.
constexpr int N = cute::size<1>(TileShape_MNK{});
static_assert(N == 64 || N == 128 || N == 192 || N == 256, "Invalid TileShape_N.");
I'm fine-tuning my gemm library with a large range of N values that fits best for fully occupying SMs. So this restriction comes short for this end.
Plus, is it a good practice to place a weight into B and an activation into A in a FC layer?
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 sm100_make_blockscaled_1sm_trivial_tiled_mma() in the CUTLASS 3.9 source and inspect the TileShape_N static assertion alongside the mxf8f6f4 1SM and 2SM constraints described in the issue. Determine whether wider N support and the A/B weight-activation arrangement are intended changes; done would require a maintainer-backed explanation or a clearly scoped design for altering the restriction.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- hpc, performance
- Issue type
- Feature
- Difficulty
- 5/5
- Estimated time
- Over a week
- Activity status
- Stale
- Clarity
- Needs clarification
- Newbie friendliness
- 25/100