[llvm] [NVVM][NVPTX] Add support for tcgen05.mma (PR #151949)

[via llvm-commits]

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 HEAD~1 HEAD --extensions h,cpp -- llvm/include/llvm/IR/NVVMIntrinsicUtils.h llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp llvm/lib/Target/NVPTX/NVPTXISelLowering.h llvm/lib/Target/NVPTX/NVPTXSubtarget.h
``````````

</details>

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

``````````diff
diff --git a/llvm/include/llvm/IR/NVVMIntrinsicUtils.h b/llvm/include/llvm/IR/NVVMIntrinsicUtils.h
index 2b041262e..f18e55c74 100644
--- a/llvm/include/llvm/IR/NVVMIntrinsicUtils.h
+++ b/llvm/include/llvm/IR/NVVMIntrinsicUtils.h
@@ -47,12 +47,7 @@ enum class CTAGroupKind : uint8_t {
   CG_2 = 2,    // cta_group::2 modifier
 };
 
-enum class Tcgen05MMAKind : uint8_t {
-  F16 = 0,
-  TF32 = 1,
-  F8F6F4 = 2,
-  I8 = 3
-};
+enum class Tcgen05MMAKind : uint8_t { F16 = 0, TF32 = 1, F8F6F4 = 2, I8 = 3 };
 
 enum class Tcgen05CollectorUsageOp : uint8_t {
   DISCARD = 0,
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index a8dd6b4da..b9d6623c8 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -1171,28 +1171,50 @@ const char *NVPTXTargetLowering::getTargetNodeName(unsigned Opcode) const {
     MAKE_CASE(NVPTXISD::TCGEN05_MMA_TENSOR_TF32_SCALE_D_DISABLE_OUTPUT_LANE_CG2)
     MAKE_CASE(NVPTXISD::TCGEN05_MMA_TENSOR_DISABLE_OUTPUT_LANE_CG1_ASHIFT)
     MAKE_CASE(NVPTXISD::TCGEN05_MMA_TENSOR_DISABLE_OUTPUT_LANE_CG2_ASHIFT)
-    MAKE_CASE(NVPTXISD::TCGEN05_MMA_TENSOR_F16_SCALE_D_DISABLE_OUTPUT_LANE_CG1_ASHIFT)
-    MAKE_CASE(NVPTXISD::TCGEN05_MMA_TENSOR_F16_SCALE_D_DISABLE_OUTPUT_LANE_CG2_ASHIFT)
-    MAKE_CASE(NVPTXISD::TCGEN05_MMA_TENSOR_TF32_SCALE_D_DISABLE_OUTPUT_LANE_CG1_ASHIFT)
-    MAKE_CASE(NVPTXISD::TCGEN05_MMA_TENSOR_TF32_SCALE_D_DISABLE_OUTPUT_LANE_CG2_ASHIFT)
+    MAKE_CASE(
+        NVPTXISD::TCGEN05_MMA_TENSOR_F16_SCALE_D_DISABLE_OUTPUT_LANE_CG1_ASHIFT)
+    MAKE_CASE(
+        NVPTXISD::TCGEN05_MMA_TENSOR_F16_SCALE_D_DISABLE_OUTPUT_LANE_CG2_ASHIFT)
+    MAKE_CASE(
+        NVPTXISD::
+            TCGEN05_MMA_TENSOR_TF32_SCALE_D_DISABLE_OUTPUT_LANE_CG1_ASHIFT)
+    MAKE_CASE(
+        NVPTXISD::
+            TCGEN05_MMA_TENSOR_TF32_SCALE_D_DISABLE_OUTPUT_LANE_CG2_ASHIFT)
     MAKE_CASE(NVPTXISD::TCGEN05_MMA_SP_SHARED_DISABLE_OUTPUT_LANE_CG1)
     MAKE_CASE(NVPTXISD::TCGEN05_MMA_SP_SHARED_DISABLE_OUTPUT_LANE_CG2)
-    MAKE_CASE(NVPTXISD::TCGEN05_MMA_SP_SHARED_F16_SCALE_D_DISABLE_OUTPUT_LANE_CG1)
-    MAKE_CASE(NVPTXISD::TCGEN05_MMA_SP_SHARED_F16_SCALE_D_DISABLE_OUTPUT_LANE_CG2)
-    MAKE_CASE(NVPTXISD::TCGEN05_MMA_SP_SHARED_TF32_SCALE_D_DISABLE_OUTPUT_LANE_CG1)
-    MAKE_CASE(NVPTXISD::TCGEN05_MMA_SP_SHARED_TF32_SCALE_D_DISABLE_OUTPUT_LANE_CG2)
+    MAKE_CASE(
+        NVPTXISD::TCGEN05_MMA_SP_SHARED_F16_SCALE_D_DISABLE_OUTPUT_LANE_CG1)
+    MAKE_CASE(
+        NVPTXISD::TCGEN05_MMA_SP_SHARED_F16_SCALE_D_DISABLE_OUTPUT_LANE_CG2)
+    MAKE_CASE(
+        NVPTXISD::TCGEN05_MMA_SP_SHARED_TF32_SCALE_D_DISABLE_OUTPUT_LANE_CG1)
+    MAKE_CASE(
+        NVPTXISD::TCGEN05_MMA_SP_SHARED_TF32_SCALE_D_DISABLE_OUTPUT_LANE_CG2)
     MAKE_CASE(NVPTXISD::TCGEN05_MMA_SP_TENSOR_DISABLE_OUTPUT_LANE_CG1)
     MAKE_CASE(NVPTXISD::TCGEN05_MMA_SP_TENSOR_DISABLE_OUTPUT_LANE_CG2)
     MAKE_CASE(NVPTXISD::TCGEN05_MMA_SP_TENSOR_DISABLE_OUTPUT_LANE_CG1_ASHIFT)
     MAKE_CASE(NVPTXISD::TCGEN05_MMA_SP_TENSOR_DISABLE_OUTPUT_LANE_CG2_ASHIFT)
-    MAKE_CASE(NVPTXISD::TCGEN05_MMA_SP_TENSOR_F16_SCALE_D_DISABLE_OUTPUT_LANE_CG1)
-    MAKE_CASE(NVPTXISD::TCGEN05_MMA_SP_TENSOR_F16_SCALE_D_DISABLE_OUTPUT_LANE_CG2)
-    MAKE_CASE(NVPTXISD::TCGEN05_MMA_SP_TENSOR_TF32_SCALE_D_DISABLE_OUTPUT_LANE_CG1)
-    MAKE_CASE(NVPTXISD::TCGEN05_MMA_SP_TENSOR_TF32_SCALE_D_DISABLE_OUTPUT_LANE_CG2)
-    MAKE_CASE(NVPTXISD::TCGEN05_MMA_SP_TENSOR_F16_SCALE_D_DISABLE_OUTPUT_LANE_CG1_ASHIFT)
-    MAKE_CASE(NVPTXISD::TCGEN05_MMA_SP_TENSOR_F16_SCALE_D_DISABLE_OUTPUT_LANE_CG2_ASHIFT)
-    MAKE_CASE(NVPTXISD::TCGEN05_MMA_SP_TENSOR_TF32_SCALE_D_DISABLE_OUTPUT_LANE_CG1_ASHIFT)
-    MAKE_CASE(NVPTXISD::TCGEN05_MMA_SP_TENSOR_TF32_SCALE_D_DISABLE_OUTPUT_LANE_CG2_ASHIFT)
+    MAKE_CASE(
+        NVPTXISD::TCGEN05_MMA_SP_TENSOR_F16_SCALE_D_DISABLE_OUTPUT_LANE_CG1)
+    MAKE_CASE(
+        NVPTXISD::TCGEN05_MMA_SP_TENSOR_F16_SCALE_D_DISABLE_OUTPUT_LANE_CG2)
+    MAKE_CASE(
+        NVPTXISD::TCGEN05_MMA_SP_TENSOR_TF32_SCALE_D_DISABLE_OUTPUT_LANE_CG1)
+    MAKE_CASE(
+        NVPTXISD::TCGEN05_MMA_SP_TENSOR_TF32_SCALE_D_DISABLE_OUTPUT_LANE_CG2)
+    MAKE_CASE(
+        NVPTXISD::
+            TCGEN05_MMA_SP_TENSOR_F16_SCALE_D_DISABLE_OUTPUT_LANE_CG1_ASHIFT)
+    MAKE_CASE(
+        NVPTXISD::
+            TCGEN05_MMA_SP_TENSOR_F16_SCALE_D_DISABLE_OUTPUT_LANE_CG2_ASHIFT)
+    MAKE_CASE(
+        NVPTXISD::
+            TCGEN05_MMA_SP_TENSOR_TF32_SCALE_D_DISABLE_OUTPUT_LANE_CG1_ASHIFT)
+    MAKE_CASE(
+        NVPTXISD::
+            TCGEN05_MMA_SP_TENSOR_TF32_SCALE_D_DISABLE_OUTPUT_LANE_CG2_ASHIFT)
   }
   return nullptr;
 

``````````

</details>


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