[llvm] [LLVM][NVPTX] Add support for ldmatrix extensions introduced in PTX 8.6 (PR #124899)

Artem Belevich via llvm-commits llvm-commits at lists.llvm.org
Thu Feb 13 16:18:12 PST 2025


================
@@ -7194,7 +7197,27 @@ class WMMA_REGINFO<WMMA_REGS r, string op>
 
     !and(!eq(op,"ldmatrix"),
          !eq(ptx_elt_type,"b16"),
-         !eq(geom, "m8n8")) : [hasSM<75>, hasPTX<65>]);
+         !eq(geom, "m8n8")) : [hasSM<75>, hasPTX<65>],
+
+    !and(!eq(op,"ldmatrix"),
+         !eq(ptx_elt_type,"b8"),
+         !eq(geom, "m16n16")) : [hasSM<100>, hasAAFeatures, hasPTX<86>],
----------------
Artem-B wrote:

Nit: We already have `hasSM100a`:
https://github.com/llvm/llvm-project/blob/3c2ba68915b268fd3b7d39bf62e19199b2cb8995/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td#L175-L180

Future cleanup opportunity, in a separate patch:

`hasAAFeatures` still rubs me the wrong way.  When I see `AA` my brain jumps straight to "what does it have to do with the Alias Analysys?". Also, it looks like 'accelerated' SM variants are here to stay.

Perhaps we can generalize `hasSM` to accept an optional "accelerated" argument, or just create `hasSMA<xx>`. 


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


More information about the llvm-commits mailing list