[QST] FP4 Tensor Core mma.sync Instruction Unsupported on SM_101 Architecture
Nobody has claimed this yet.
- Dominant language
- C++
- Stars
- 10.5k
- Forks
- 2.1k
- Avg merge
- 3d 11h
- Merged PRs (30d)
- 7
Description
The current kernel uses the following FP4 Tensor Core instruction:
mma.sync.aligned.m16n8k64.row.col.kind::mxf4nvf4.block_scale.scale_vec::4X.f32.e2m1.e2m1.f32.ue4m3
This instruction is only supported on SM_120+ (Hopper/Blackwell architectures). Our target platform is SM_101, which does not have native FP4 Tensor Core support, so the instruction cannot execute.
Is there a recommended method to emulate this FP4 Tensor Core MMA on SM_101 while maintaining equivalent numerical results?
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 by tracing the kernel that emits mma.sync.aligned.m16n8k64.row.col.kind::mxf4nvf4.block_scale.scale_vec::4X.f32.e2m1.e2m1.f32.ue4m3 and checking the SM_101 and SM_120 capability paths. No files or tests are named in the issue, so identify the relevant CUTLASS kernel and validation coverage first. Done means an SM_101-compatible emulation exists and produces equivalent numerical results.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- backend
- Issue type
- Feature
- Difficulty
- 5/5
- Estimated time
- Over a week
- Activity status
- Quiet
- Clarity
- Needs clarification
- Newbie friendliness
- 25/100