NVIDIA / NVIDIA/cutlass

[QST][CuteDSL] warp mma support

Open
#2,825 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

I have noticed that CuteDSL only supports fp16/bf16 warp mma with shape m16n8k16 and m16n8k8 now.

https://github.com/NVIDIA/cutlass/blob/ec8daf642d69fc31352ac6fa6e14a0de9019604b/python/CuTeDSL/cutlass/cute/nvgpu/warp/mma.py

Are there any other support plans in the future, such as

  • Turing support
  • Volta support (with mma shape m8n8k4)
  • B1 / INT4 / INT8 / FP4 / FP6 / FP8 / TF32 support

BTW, the documentation comments for MmaF16BF16Op seem to be incorrect. My understanding is that this is not tcgen05, right?

https://github.com/NVIDIA/cutlass/blob/ec8daf642d69fc31352ac6fa6e14a0de9019604b/python/CuTeDSL/cutlass/cute/nvgpu/warp/mma.py#L43-L50

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 by reading python/CuTeDSL/cutlass/cute/nvgpu/warp/mma.py, especially the supported warp MMA shapes and the MmaF16BF16Op documentation comments. Establish which architecture and data-type combinations are intended for support and verify the tcgen05 documentation claim with maintainers. Done would mean a decided support scope and accurate documentation.

Written by the indexing model from the issue text.

Assessment

Tech stack
python
Domain
hpc
Issue type
Feature
Difficulty
5/5
Estimated time
Over a week
Activity status
Quiet
Clarity
Needs clarification
Newbie friendliness
30/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.