[llvm] [NVPTX] Misc table-gen cleanup (NFC) (PR #142877)

[Artem Belevich via llvm-commits]

================
@@ -7042,36 +5668,33 @@ def tcgen05_wait_st: BasicNVPTXInst<(outs), (ins), "tcgen05.wait::st.sync.aligne
   [(int_nvvm_tcgen05_wait_st)]>,
   Requires<[hasTcgen05Instructions]>;
 
-multiclass TCGEN05_COMMIT_INTR<NVPTXRegClass rc, string AS, string num> {
-  defvar prefix = "tcgen05.commit.cta_group::" # num;
-  defvar suffix = ".mbarrier::arrive::one.shared::cluster";
+multiclass TCGEN05_COMMIT_INTR<string AS, string num> {
+  defvar prefix = "tcgen05.commit.cta_group::" # num #".mbarrier::arrive::one.shared::cluster";
 
   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);
 
-  def NAME : NVPTXInst<(outs), (ins rc:$mbar),
-             !strconcat(prefix, suffix, ".b64 [$mbar];"),
-             [(Intr rc:$mbar)]>,
+  def "" : BasicNVPTXInst<(outs), (ins ADDR:$mbar),
+             prefix # ".b64",
+             [(Intr addr:$mbar)]>,
              Requires<[hasTcgen05Instructions]>;
-  def NAME # _MC : NVPTXInst<(outs), (ins rc:$mbar, Int16Regs:$mc),
-                   !strconcat(prefix, suffix, ".multicast::cluster.b64 [$mbar], $mc;"),
-                   [(IntrMC rc:$mbar, Int16Regs:$mc)]>,
+  def _MC : BasicNVPTXInst<(outs), (ins ADDR:$mbar, Int16Regs:$mc),
+                   prefix # ".multicast::cluster.b64",
+                   [(IntrMC addr:$mbar, Int16Regs:$mc)]>,
                    Requires<[hasTcgen05Instructions]>;
 }
 
-defm TCGEN05_COMMIT_CG1 : TCGEN05_COMMIT_INTR<Int64Regs, "", "1">;
-defm TCGEN05_COMMIT_CG2 : TCGEN05_COMMIT_INTR<Int64Regs, "", "2">;
-defm TCGEN05_COMMIT_S64_CG1 : TCGEN05_COMMIT_INTR<Int64Regs, "shared", "1">;
-defm TCGEN05_COMMIT_S64_CG2 : TCGEN05_COMMIT_INTR<Int64Regs, "shared", "2">;
-defm TCGEN05_COMMIT_S32_CG1 : TCGEN05_COMMIT_INTR<Int32Regs, "shared", "1">;
-defm TCGEN05_COMMIT_S32_CG2 : TCGEN05_COMMIT_INTR<Int32Regs, "shared", "2">;
----------------
Artem-B wrote:

Were S32 variants removed intentionally?

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