[llvm] [NVPTX] Add sparse MMA intrinsics (PR #150950)

[Durgadoss R via llvm-commits]

================
@@ -170,14 +170,62 @@ class StrJoin<string sep, list<string> str_list> {
 // Geom: m<M>n<N>k<K>. E.g. m8n32k16
 // Frag: [a|b|c|d] ([x1|x2|x4] for ldmatrix)
 // PtxEltType: PTX type for the element.
-class WMMA_REGS<string Geom, string Frag, string PtxEltType> {
+class WMMA_REGS<string Geom, string Frag, string PtxEltType, bit IsSparse = false> {
   string geom = Geom;
   string frag = Frag;
   string ptx_elt_type = PtxEltType;
   string gft = Geom#":"#Frag#":"#ptx_elt_type;
   string gf = Geom#":"#Frag;
   string ft = frag#":"#ptx_elt_type;
   list<LLVMType> regs = !cond(
+    // mma sparse ops use other fragments for some arguments
+    !and(!eq(gft, "m16n8k16:a:bf16"), !eq(IsSparse, true)) : !listsplat(llvm_i32_ty, 2),
+    !and(!eq(gft, "m16n8k16:a:f16"), !eq(IsSparse, true)) : !listsplat(llvm_v2f16_ty, 2),
----------------
durga4github wrote:

As long as we have thought through that, we are good.

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