[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:13:03 PDT 2026
https://github.com/weimin023 updated https://github.com/llvm/llvm-project/pull/211841
>From e2100d130383813da0fc7f3f45457f7c51dbfb29 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 b15ee81291e5c..30a87ec7ff62c 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -3588,6 +3588,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 |
|----------|-----------|---------|---------|----------|----------|-------------------|
@@ -3607,6 +3609,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 15d704e9ca0aa48666038373ae3eed740f099fad 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 30a87ec7ff62c..248a864be3ab4 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -3591,26 +3591,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