如何理解以下CUDA内联汇编代码片段中的参数？
......
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呢？
好吧，那些矩阵运算PTX指令确实很难理解！
但是，让我们仔细阅读PTX ISA参考资料的相关部分：
- 被乘数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。