[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