[Mlir-commits] [mlir] [mlir][doc] Document FP8 e4m3/e5m2 mma.sync formats (PR #211841)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Sun Oct 11 00:08:54 PDT 2026
https://github.com/weimin023 updated https://github.com/llvm/llvm-project/pull/211841
>From cadf5bedd6d3286a073085ed21a66bf538590c16 Mon Sep 17 00:00:00 2001
From: weimin023 <tnwilly at gmail.com>
Date: Fri, 24 Jul 2026 15:49:05 +0000
Subject: [PATCH 1/2] [doc] document FP8 mma.sync
Signed-off-by: weimin023 <tnwilly at gmail.com>
---
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 4 ++++
1 file changed, 4 insertions(+)
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 1f76218936e16..e2ae8ccfb0aef 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -3396,6 +3396,8 @@ def NVVM_MmaOp : NVVM_Op<"mma.sync", [AttrSizedOperandSegments]> {
`#nvvm.mma_layout<row>` and `#nvvm.mma_layout<col>` respectively, but other
combinations are possible for certain layouts according to the table below.
+ For FP8 operations, A and B can independently use `e4m3` or `e5m2`.
+
```
| A/B Type | Shape | ALayout | BLayout | A Type | B Type | C/D Type |
|----------|-----------|---------|---------|----------|----------|-------------------|
@@ -3415,6 +3417,8 @@ def NVVM_MmaOp : NVVM_Op<"mma.sync", [AttrSizedOperandSegments]> {
| | m16n8k64 | row | col | 4x i32 | 2x i32 | 4x i32 |
| b1 | m8n8k128 | row | col | 1x i32 | 1x i32 | 2x i32 |
| | m16n8k128 | row | col | 2x i32 | 1x i32 | 4x i32 |
+ | e4m3/e5m2 | .m16n8k16 | row | col | 2x i32 | 1x i32 | 2x f16x2 or 4 f32 |
+ | | .m16n8k32 | row | col | 4x i32 | 2x i32 | 2x f16x2 or 4 f32 |
```
>From 7700aa8ee5628fc11340bd75ef57680ce0d17085 Mon Sep 17 00:00:00 2001
From: weimin023 <tnwilly at gmail.com>
Date: Sun, 11 Oct 2026 07:08:30 +0000
Subject: [PATCH 2/2] Fix FP8 mma docs
---
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 40 +++++++++++----------
1 file changed, 21 insertions(+), 19 deletions(-)
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index e2ae8ccfb0aef..d234693883dde 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -3399,26 +3399,28 @@ def NVVM_MmaOp : NVVM_Op<"mma.sync", [AttrSizedOperandSegments]> {
For FP8 operations, A and B can independently use `e4m3` or `e5m2`.
```
- | A/B Type | Shape | ALayout | BLayout | A Type | B Type | C/D Type |
- |----------|-----------|---------|---------|----------|----------|-------------------|
- | f64 | .m8n8k4 | row | col | 1x f64 | 1x f64 | 2x f64 |
- | f16 | .m8n8k4 | row/col | row/col | 2x f16x2 | 2x f16x2 | 4x f16x2 or 8xf32 |
- | | .m16n8k8 | row | col | 2x f16x2 | 1x f16x2 | 2x f16x2 or 4 f32 |
- | | .m16n8k16 | row | col | 4x f16x2 | 2x f16x2 | 2x f16x2 or 4 f32 |
- | bf16 | .m16n8k8 | row | col | 2x i32 | 1x i32 | 4x f32 |
- | | .m16n8k16 | row | col | 4x i32 | 2x i32 | 4x f32 |
- | tf32 | .m16n8k4 | row | col | 2x i32 | 1x i32 | 4x f32 |
- | | .m16n8k8 | row | col | 4x i32 | 2x i32 | 2x f16x2 or 4 f32 |
- | u8/s8 | .m8n8k16 | row | col | 1x i32 | 1x i32 | 2x i32 |
- | | .m16n8k16 | row | col | 2x i32 | 1x i32 | 4x i32 |
- | | .m16n8k32 | row | col | 4x i32 | 2x i32 | 4x i32 |
- | u4/s4 | .m8n8k32 | row | col | 1x i32 | 1x i32 | 2x i32 |
- | | m16n8k32 | row | col | 2x i32 | 1x i32 | 4x i32 |
- | | m16n8k64 | row | col | 4x i32 | 2x i32 | 4x i32 |
- | b1 | m8n8k128 | row | col | 1x i32 | 1x i32 | 2x i32 |
- | | m16n8k128 | row | col | 2x i32 | 1x i32 | 4x i32 |
+ | A/B Type | Shape | ALayout | BLayout | A Type | B Type | C/D Type |
+ |-----------|-----------|---------|---------|----------|----------|-------------------|
+ | f64 | .m8n8k4 | row | col | 1x f64 | 1x f64 | 2x f64 |
+ | f16 | .m8n8k4 | row/col | row/col | 2x f16x2 | 2x f16x2 | 4x f16x2 or 8xf32 |
+ | | .m16n8k8 | row | col | 2x f16x2 | 1x f16x2 | 2x f16x2 or 4 f32 |
+ | | .m16n8k16 | row | col | 4x f16x2 | 2x f16x2 | 2x f16x2 or 4 f32 |
+ | bf16 | .m16n8k8 | row | col | 2x i32 | 1x i32 | 4x f32 |
+ | | .m16n8k16 | row | col | 4x i32 | 2x i32 | 4x f32 |
+ | tf32 | .m16n8k4 | row | col | 2x i32 | 1x i32 | 4x f32 |
+ | | .m16n8k8 | row | col | 4x i32 | 2x i32 | 2x f16x2 or 4 f32 |
+ | u8/s8 | .m8n8k16 | row | col | 1x i32 | 1x i32 | 2x i32 |
+ | | .m16n8k16 | row | col | 2x i32 | 1x i32 | 4x i32 |
+ | | .m16n8k32 | row | col | 4x i32 | 2x i32 | 4x i32 |
+ | u4/s4 | .m8n8k32 | row | col | 1x i32 | 1x i32 | 2x i32 |
+ | | m16n8k32 | row | col | 2x i32 | 1x i32 | 4x i32 |
+ | | m16n8k64 | row | col | 4x i32 | 2x i32 | 4x i32 |
+ | b1 | m8n8k128 | row | col | 1x i32 | 1x i32 | 2x i32 |
+ | | m16n8k128 | row | col | 2x i32 | 1x i32 | 4x i32 |
| e4m3/e5m2 | .m16n8k16 | row | col | 2x i32 | 1x i32 | 2x f16x2 or 4 f32 |
- | | .m16n8k32 | row | col | 4x i32 | 2x i32 | 2x f16x2 or 4 f32 |
+ | | .m16n8k32 | row | col | 4x i32 | 2x i32 | 2x f16x2 or 4 f32 |
+
+ [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-shape)
```
More information about the Mlir-commits
mailing list