理解PTX指令mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32的参数

【问题标题】：Understanding the parameters of PTX instruction mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32理解PTX指令mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32的参数
【发布时间】：2022-01-13 04:30:18
【问题描述】：

如何理解CUDA内联汇编代码下面sn-p中的参数？

......
        asm volatile( \
            "mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 \n" \
            "    {%0, %1, %2, %3}, \n" \
            "    {%4, %5, %6, %7}, \n" \
            "    {%8, %9}, \n" \
            "    {%0, %1, %2, %3}; \n" \
                    : "+f"(  elt(0)), "+f"(  elt(1)), "+f"(  elt(2)), "+f"(  elt(3))
                    :  "r"(a.reg(0)),  "r"(a.reg(1)),  "r"(a.reg(2)),  "r"(a.reg(3))
                    ,  "r"(b.reg(0)),  "r"(b.reg(1)));
......

D(16x8) = A(16x16) * B(16x8) + C(16*8) 的矩阵乘法和加法。那么为什么 D:A:B:C 的参数计数是 4:4:2:4 而不是 2:4:2:2？

【问题讨论】：

标签： cuda matrix-multiplication inline-assembly ptx

【解决方案1】：

嗯，那些矩阵运算 PTX 指令肯定是相当多的绕圈子！

但是-让我们仔细阅读PTX ISA reference的relevant section：

被乘数 A，元素类型 .f16：包含四个 .f16x2 寄存器的向量表达式。
被乘数 B，元素类型 .f16：包含两个 .f16x2 寄存器的向量表达式。
加数 C，元素类型 .f32：包含四个 .f32 寄存器的向量表达式。
结果 D，元素类型 .f32：包含四个 .f32 寄存器的向量表达式。

A x B 将矩阵大小减小到 B 的大小，向量表达式大小为 2 而不是 4，因此您可能期望 C 的向量表达式大小为 2，h-o-w-e-v-e-r - 您增加了元素大小到.f32，因此你有一个.f32，而不是每个.f16x2，所以你保留了一个4元素向量表达式而不是2来解释这个事实。

因此，D:A:B:C 为 4:4:2:4。

【讨论】：