NVIDIA / NVIDIA/cccl

[FEA]: Add cuda::ptx wrappers for warp-level mma.sync instructions

Open
#8,146 6 comments 0 reactions 0 assignees View on GitHub
needs triage
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

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.