[llvm] [NFC][NVVM][NVPTX] Moved common code for tcgen05.mma to the base class (PR #176327)

[Pradeep Kumar via llvm-commits]

================
@@ -6071,195 +6060,142 @@ foreach sp = [0, 1] in {
                                           !eq(kind, "tf32")), [0, 1], [0]) in {
             foreach ashift = !if(!eq(space, "tensor"), [0, 1], [0]) in {
               def :
-                Tcgen05MMADisableOutputLaneInst<sp, space, kind, cta_group,
-                                               collector_usage, scale_input_d,
-                                               ashift>;
-            }
-          }
-        }
-      }
-    }
-  }
-}
+                Tcgen05MMADisableOutputLaneInst<sparse, space, kind, cta_group,
+                                                collector_usage, scale_input_d,
+                                                ashift>;
+            } // ashift
+          } // scale_input_d
+        } // collector_usage
+      } // cta_group
+    } // kind
+  } // space
+} // sparse
 
-class Tcgen05MMABlockScaleInst<bit Sp, string ASpace, string KindStr,
-                     int CtaGroup, string ScaleVecSize, string CollectorUsageStr>:
-         NVPTXInst<(outs), (ins), "?", []>, Requires<[]> {
+//
+// tcgen05.mma.block_scale Instructions
+//
+
+class Tcgen05MMABlockScaleInst<bit IsSparse, string ASpace, string Kind,
+                               int CtaGroup, string ScaleVecSize,
+                               string CollectorUsage>:
+        Tcgen05MMABase<IsSparse, ASpace, Kind, CtaGroup, CollectorUsage> {
 
   let Predicates = !cond(
-    !and(!eq(Sp, 1),
-         !eq(KindStr, "mxf4")) : [callSubtarget<"hasTcgen05MMASparseMxf4">],
-    !and(!eq(Sp, 1),
-         !eq(KindStr, "mxf4nvf4")) : [callSubtarget<"hasTcgen05MMASparseMxf4nvf4">],
+    !and(IsSparse,
+         !eq(Kind, "mxf4")) : [callSubtarget<"hasTcgen05MMASparseMxf4">],
+    !and(IsSparse,
+         !eq(Kind, "mxf4nvf4")) : [callSubtarget<"hasTcgen05MMASparseMxf4nvf4">],
     !ne(ScaleVecSize, "") : [callSubtarget<"hasTcgen05InstSupport">, hasPTX<88>],
     true : [callSubtarget<"hasTcgen05InstSupport">]
   );
 
   Intrinsic Intrin = !cast<Intrinsic>(
-                             NVVM_TCGEN05_MMA_BLOCKSCALE<Sp, ASpace, KindStr, ScaleVecSize>.record_name);
-
-  dag SparseMetadataIns = !if(!eq(Sp, 1), (ins B32:$spmetadata), (ins));
-  dag SparseMetadataIntr = !if(!eq(Sp, 1), (Intrin i32:$spmetadata), (Intrin));
-  string SparseMetadataStr = !if(!eq(Sp, 1), ", [$spmetadata]", "");
+                             NVVM_TCGEN05_MMA_BLOCKSCALE<IsSparse, ASpace, Kind, ScaleVecSize>.record_name);
 
-  int KindVal = !cond(
-                  !eq(KindStr, "mxf8f6f4") : 0,
-                  !eq(KindStr, "mxf4")     : 1,
-                  !eq(KindStr, "mxf4nvf4") : 2,
+  let KindVal = !cond(
----------------
schwarzschild-radius wrote:

I think we can remove this as it is not used?

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