[FEA]: Add cuda::ptx wrappers for warp-level mma.sync instructions
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 486
- Avg merge
- 2d 6h
- Merged PRs (30d)
- 295
Description
### Is this a duplicate?
- [x] I confirmed there appear to be no [duplicate issues](https://github.com/NVIDIA/cccl/issues) for this request and that I agree to the [Code of Conduct](CODE_OF_CONDUCT.md)
### Area
libcu++
### Is your feature request related to a problem? Please describe.
Hi,
I am working on a fused FP4 attention kernel targeting SM120 (consumer Blackwell, RTX 5070 Ti). On SM120, tcgen05.mma is not available, so the only path to FP4 Tensor Cores is the warp-level mma.sync instruction:
`mma.sync.aligned.kind::mxf8f6f4.block_scale.scale_vec::1X.m16n8k32.row.col.f32.e2m1.e2m1.f32.ue8m0`
Looking at the instruction table in the cuda::ptx documentation, all warp-level MMA instructions (mma, ldmatrix, stmatrix, wmma.mma, mma.sp, movmatrix) are listed as "No" under "Available in libcu++".
Currently the only way to use this instruction is through asm volatile inline PTX, which works but loses the type safety and compile-time validation that cuda::ptx provides for other instructions.
Happy to help test if a wrapper is in progress.
Reference: https://nvidia.github.io/cccl/unstable/libcudacxx/ptx/instructions.html
### Describe the solution you'd like
Would the team consider adding wrappers for warp-level mma.sync, at least for the block-scaled FP4/FP6/FP8 variants (kind::mxf8f6f4)? These are the only MMA instructions available on SM120 consumer GPUs, and I suspect anyone doing low-level Tensor Core work on the 50-series will run into the same gap.
### Describe alternatives you've considered
_No response_
### Additional context
Discussed with Federico Busato
Contributor guide
Assessment
This issue has not been assessed yet.