[all-commits] [llvm/llvm-project] 69ee6b: [NVVM] Add FP8 (e4m3/e5m2) support to dense mma.sy...

[WMC via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: 69ee6b114ac3d3c6652b39d0d57b292bcfa3a81c
      https://github.com/llvm/llvm-project/commit/69ee6b114ac3d3c6652b39d0d57b292bcfa3a81c
  Author: WMC <tnwilly at gmail.com>
  Date:   2026-07-14 (Tue, 14 Jul 2026)

  Changed paths:
    M mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
    M mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
    M mlir/test/Dialect/LLVMIR/nvvm.mlir
    M mlir/test/Target/LLVMIR/nvvmir.mlir

  Log Message:
  -----------
  [NVVM] Add FP8 (e4m3/e5m2) support to dense mma.sync (#207307)

## Motivation
I'm adding NVIDIA FP8 mma.sync support in IREE for sm_89/sm_120, and
IREE's codegen lowers through `nvgpu.mma.sync` on this path, which in
turn lowers to `nvvm.mma.sync`. This MLIR-level gap is blocking that
work, so I'm upstreaming the fix here first. This is split into two PRs
per review feedback — this one covers `nvvm.mma.sync` only; a follow-up
PR will add the corresponding `nvgpu.mma.sync` support built on top of
this.

## Description:

`nvvm.mma.sync` (warp-level, dense) don't support FP8 (e4m3/e5m2)
operands, even though the underlying PTX/NVPTX intrinsics have supported
FP8 `mma.sync.m16n8k32` since sm_89. FP8 currently only exists on the
sparse mma path and on warpgroup-level `wgmma`; the warp-level dense
path was never updated.

This adds FP8 (e4m3/e5m2) support to `nvvm.mma.sync`, covering
`m16n8k16` and `m16n8k32`.

## Verification
`$ mlir-opt mlir/test/Dialect/LLVMIR/nvvm.mlir -split-input-file
-verify-diagnostics | FileCheck mlir/test/Dialect/LLVMIR/nvvm.mlir
`

```
%0 = nvvm.mma.sync A[...] B[...] C[...]
     {multiplicandAPtxType = #nvvm.mma_type<e4m3>,
      multiplicandBPtxType = #nvvm.mma_type<e4m3>,
      shape = #nvvm.shape<m = 16, n = 8, k = 32>}
     : (i32, i32, f32) -> !llvm.struct<(f32, f32, f32, f32)>

```

`$ mlir-translate -mlir-to-llvmir mlir/test/Target/LLVMIR/nvvmir.mlir
-split-input-file -verify-diagnostics | FileCheck
mlir/test/Target/LLVMIR/nvvmir.mlir
`

```
%8 = call { float, float, float, float } @llvm.nvvm.mma.m16n8k16.row.col.f32.e4m3.e4m3.f32(i32 %0, i32 %1, i32 %2, float %3, float %4, float %5, float %6)
%11 = call { float, float, float, float } @llvm.nvvm.mma.m16n8k32.row.col.f32.e4m3.e5m2.f32(i32 %0, i32 %1, i32 %2, i32 %3, i32 %4, i32 %5, float %6, float %7, float %8, float %9)

```

All existing tests in nvvm.mlir (including -verify-diagnostics error
cases) still pass — no regression.

Co-authored with Claude Sonnet 5

Signed-off-by: weimin023 <tnwilly at gmail.com>



To unsubscribe from these emails, change your notification settings at https://github.com/llvm/llvm-project/settings/notifications