[Mlir-commits] [mlir] [mlir][doc] Document FP8 e4m3/e5m2 mma.sync formats (PR #211841)

[llvmlistbot at llvm.org]

https://github.com/weimin023 created https://github.com/llvm/llvm-project/pull/211841

## Description

Update the documentation for the supported FP8 `mma.sync` formats introduced in #207307

>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] [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 |
     ```