[Mlir-commits] [llvm] [mlir] [NVPTX] Add Rubin extensions to tcgen05.commit (PR #211577)

[Durgadoss R]

================
@@ -5711,25 +5711,41 @@ def tcgen05_wait_ld: NullaryInst<"tcgen05.wait::ld.sync.aligned", int_nvvm_tcgen
 def tcgen05_wait_st: NullaryInst<"tcgen05.wait::st.sync.aligned", int_nvvm_tcgen05_wait_st>;
 
 multiclass TCGEN05_COMMIT_INTR<string AS, string num> {
-  defvar prefix = "tcgen05.commit.cta_group::" # num #".mbarrier::arrive::one.shared::cluster";
+  defvar prefix = "tcgen05.commit.cta_group::" # num #
+                  ".mbarrier::arrive::one.shared::cluster";
+  defvar prefix_smem_a_read = prefix # ".sync_restrict::shared::read::mma::a";
 
+  defvar intr_prefix = "int_nvvm_tcgen05_commit";
   defvar intr_suffix = !if(!eq(AS, "shared"), "_shared", "") # "_cg" # num;
-  defvar Intr = !cast<Intrinsic>("int_nvvm_tcgen05_commit" # intr_suffix);
-  defvar IntrMC = !cast<Intrinsic>("int_nvvm_tcgen05_commit_mc" # intr_suffix);
+  defvar Intr = !cast<Intrinsic>(intr_prefix # intr_suffix);
+  defvar IntrMC = !cast<Intrinsic>(intr_prefix # "_mc" # intr_suffix);
+  defvar IntrSmemARead =
+      !cast<Intrinsic>(intr_prefix # "_smem_a_read" # intr_suffix);
+  defvar IntrSmemAReadMC =
+      !cast<Intrinsic>(intr_prefix # "_smem_a_read_mc" # intr_suffix);
 
   def "" : BasicNVPTXInst<(outs), (ins ADDR:$mbar),
              prefix # ".b64",
              [(Intr addr:$mbar)]>;
   def _MC : BasicNVPTXInst<(outs), (ins ADDR:$mbar, B16:$mc),
-                   prefix # ".multicast::cluster.b64",
-                   [(IntrMC addr:$mbar, B16:$mc)]>;
+              prefix # ".multicast::cluster.b64",
+              [(IntrMC addr:$mbar, i16:$mc)]>;
+  let Predicates = [callSubtarget<"hasTcgen05RubinFamilySupport">] in {
+    def _MC_32B : BasicNVPTXInst<(outs), (ins ADDR:$mbar, B32:$mc),
+                    prefix # ".multicast::cluster::32b.b64",
+                    [(IntrMC addr:$mbar, i32:$mc)]>;
+    def _SMRA : BasicNVPTXInst<(outs), (ins ADDR:$mbar),
----------------
durga4github wrote:

Can we give it a bit more explicit name?
`SMEM_RA_` or `SMEM_READ` or fully `SMEM_A_READ`..

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