[Mlir-commits] [mlir] [NVGPU] Add FP8 (e4m3/e5m2) support to nvgpu.mma.sync (PR #207342)

[llvmlistbot at llvm.org]

================
@@ -0,0 +1,12 @@
+// RUN: mlir-opt %s -convert-nvgpu-to-nvvm -split-input-file | FileCheck %s
+
+// Test that FP8 (e4m3) nvgpu.mma.sync lowers to nvvm.mma.sync with the
+// correct multiplicand PTX type.
+func.func @fp8_mma_sync(%arg0: vector<4x4xf8E4M3FN>, %arg1: vector<2x4xf8E4M3FN>, %arg2: vector<2x2xf32>) -> vector<2x2xf32> {
----------------
weimin023 wrote:

Hi @kvederni, 
Got it! I renamed the test file and added E4M3 and E5M2 coverage for both m16n8k16 and m16n8k32.

I did not add mixed E4M3/E5M2 A/B cases because `nvgpu.mma.sync` currently has an existing verifier constraint requiring matrixA and matrixB to have the same element type. 

https://github.com/llvm/llvm-project/pull/207342