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

[llvmlistbot at llvm.org]

github-actions[bot] wrote:

<!--LLVM CODE FORMAT COMMENT: {clang-format}-->


:warning: C/C++ code formatter, clang-format found issues in your code. :warning:

<details>
<summary>
You can test this locally with the following command:
</summary>

``````````bash
git-clang-format --diff origin/main HEAD --extensions cpp -- mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp --diff_from_common_commit
``````````

:warning:
The reproduction instructions above might return results for more than one PR
in a stack if you are using a stacked PR workflow. You can limit the results by
changing `origin/main` to the base branch/commit you want to compare against.
:warning:

</details>

<details>
<summary>
View the diff from clang-format here.
</summary>

``````````diff
diff --git a/mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp b/mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp
index 198fd9586..f5232283d 100644
--- a/mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp
+++ b/mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp
@@ -184,8 +184,8 @@ static LogicalResult verifyMmaSyncOp(Operation *op,
     numElementA = 1;
     numElementB = 1;
   } else if (aType.isF32() || aType.isBF16() || aType.isF16() ||
-             aType.isInteger(8) || aType.isInteger(4) ||
-             aType.isF8E4M3FN() || aType.isF8E5M2()) {
+             aType.isInteger(8) || aType.isInteger(4) || aType.isF8E4M3FN() ||
+             aType.isF8E5M2()) {
     // 8-by-8-128b fundamental tensor core tile size
     int operandBitwidth = aType.getIntOrFloatBitWidth();
     shapeK = 128 / operandBitwidth; // 128b wide shapeK

``````````

</details>


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