[Mlir-commits] [llvm] [mlir] [MLIR][NVVM] Support stmatrix intrinsics (PR #148377)

[Durgadoss R]

================
@@ -164,6 +164,50 @@ static llvm::Intrinsic::ID getLdMatrixIntrinsicId(NVVM::MMALayout layout,
   }
 }
 
+/// Return the intrinsic ID associated with stmatrix for the given paramters.
+static llvm::Intrinsic::ID
+getStMatrixIntrinsicId(NVVM::MMALayout layout, int32_t num,
+                       NVVM::LdStMatrixShapeAttr shape,
+                       NVVM::LdStMatrixEltType eltType) {
+  if (shape.getM() == 8 && shape.getN() == 8) {
+    if (eltType == NVVM::LdStMatrixEltType::B16) {
+      if (layout == NVVM::MMALayout::row) {
+        switch (num) {
+        case 1:
+          return llvm::Intrinsic::nvvm_stmatrix_sync_aligned_m8n8_x1_b16;
+        case 2:
+          return llvm::Intrinsic::nvvm_stmatrix_sync_aligned_m8n8_x2_b16;
+        case 4:
+          return llvm::Intrinsic::nvvm_stmatrix_sync_aligned_m8n8_x4_b16;
+        }
+      } else {
+        switch (num) {
+        case 1:
+          return llvm::Intrinsic::nvvm_stmatrix_sync_aligned_m8n8_x1_trans_b16;
+        case 2:
+          return llvm::Intrinsic::nvvm_stmatrix_sync_aligned_m8n8_x2_trans_b16;
+        case 4:
+          return llvm::Intrinsic::nvvm_stmatrix_sync_aligned_m8n8_x4_trans_b16;
+        }
+      }
+    }
+  } else if (shape.getM() == 16 && shape.getN() == 8) {
+    if (eltType == NVVM::LdStMatrixEltType::B8) {
+      if (layout == NVVM::MMALayout::col) {
----------------
durga4github wrote:

We are checking for these conditions in the Verifier.
So, I guess, we can remove the redundant checks here.
(i.e. keep only the switch-case statements here)

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