[Mlir-commits] [mlir] [MLIR][NVVM] Add tcgen05.mma MLIR Ops (PR #164356)

[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/Dialect/LLVMIR/IR/NVVMDialect.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/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index 475625fd1..cdb768270 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -2606,19 +2606,19 @@ Tcgen05MMAOp::getIntrinsicIDAndArgs(Operation &op, LLVM::ModuleTranslation &mt,
     static constexpr llvm::Intrinsic::ID tcgen05MMAIDs[2][2] = {
         // shared
         {llvm::Intrinsic::nvvm_tcgen05_mma_shared,
-          llvm::Intrinsic::nvvm_tcgen05_mma_shared},
-         // tensor
-         {llvm::Intrinsic::nvvm_tcgen05_mma_tensor,
-          llvm::Intrinsic::nvvm_tcgen05_mma_tensor_ashift}};
+         llvm::Intrinsic::nvvm_tcgen05_mma_shared},
+        // tensor
+        {llvm::Intrinsic::nvvm_tcgen05_mma_tensor,
+         llvm::Intrinsic::nvvm_tcgen05_mma_tensor_ashift}};
 
     // Scaled [isATensor][enableAshift]
     static constexpr llvm::Intrinsic::ID scaledIDs[2][2] = {
         // shared
         {llvm::Intrinsic::nvvm_tcgen05_mma_shared_scale_d,
-          llvm::Intrinsic::nvvm_tcgen05_mma_shared_scale_d},
-         // tensor
-         {llvm::Intrinsic::nvvm_tcgen05_mma_tensor_scale_d,
-          llvm::Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_ashift}};
+         llvm::Intrinsic::nvvm_tcgen05_mma_shared_scale_d},
+        // tensor
+        {llvm::Intrinsic::nvvm_tcgen05_mma_tensor_scale_d,
+         llvm::Intrinsic::nvvm_tcgen05_mma_tensor_scale_d_ashift}};
 
     // Scaled + disable output lane [isATensor][enableAshift][ctaGroup-1]
     static constexpr llvm::Intrinsic::ID disableOutputLaneIDs[2][2][2] = {
@@ -2754,19 +2754,19 @@ mlir::NVVM::IDArgPair Tcgen05MMASpOp::getIntrinsicIDAndArgs(
     static constexpr llvm::Intrinsic::ID tcgen05MMAIDs[2][2] = {
         // shared
         {llvm::Intrinsic::nvvm_tcgen05_mma_sp_shared,
-          llvm::Intrinsic::nvvm_tcgen05_mma_sp_shared},
-         // tensor
-         {llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor,
-          llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_ashift}};
+         llvm::Intrinsic::nvvm_tcgen05_mma_sp_shared},
+        // tensor
+        {llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor,
+         llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_ashift}};
 
     // Scaled [isATensor][enableAshift]
     static constexpr llvm::Intrinsic::ID scaledIDs[2][2] = {
         // shared
         {llvm::Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d,
-          llvm::Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d},
-         // tensor
-         {llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d,
-          llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_ashift}};
+         llvm::Intrinsic::nvvm_tcgen05_mma_sp_shared_scale_d},
+        // tensor
+        {llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d,
+         llvm::Intrinsic::nvvm_tcgen05_mma_sp_tensor_scale_d_ashift}};
 
     // Scaled + disable output lane [isATensor][enableAshift][ctaGroup-1]
     static constexpr llvm::Intrinsic::ID disableOutputLaneIDs[2][2][2] = {

``````````

</details>


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