[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


================
@@ -536,13 +556,18 @@ class NVVM_MMA_SUPPORTED<list<WMMA_REGS> frags, string layout_a, string layout_b
 // if NVVM_LDMATRIX_SUPPORTED<...>.ret then
 //   def : FOO<>; // The record will only be defined for supported ops.
 //
-class NVVM_LDMATRIX_SUPPORTED<WMMA_REGS frag> {
+class NVVM_LDMATRIX_SUPPORTED<WMMA_REGS frag, bit trans> {
   string g = frag.geom;
   string t = frag.ptx_elt_type;
 
   bit ret = !cond(
-    // Only currently support m8n8 and b16
     !and(!eq(g, "m8n8"), !eq(t, "b16")): true,
+    !and(!eq(g, "m16n16"), !eq(t, "b8"), !eq(trans, 1)): true,
----------------
Artem-B wrote:

Do I understand it correctly that for `m8n8.b16` we do intend to generate both transposed and non-transposed variants?
Patch description only mentions m16n16 and m8n16, so I'm not sure if unchanged m8n8 condition is intentional. 


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


More information about the llvm-commits mailing list