Skip to main content

DCU 上的 Matrix Core 编程（part2）

· 5 min read
genshen
genshen
student in CS

Matrix Core 上矩阵乘指令的基本格式​

我们考虑矩阵乘：

D=A×B+CD = A \times B + C

其中，A 矩阵是 M×KM \times K 的（M 行、K列，后续记号也是一样的含义）， B 矩阵是 K×NK \times N的， C、D 矩阵都是 M×NM \times N 的。

对应的 API 的格式为： d = __builtin_amdgcn_mmac_CDFmt_MxNxKABFmt(a, b, c); 其中，约定如下：

  • ABFmt：AA、BB 矩阵中元素的精度，可以是 f16、f32、f64、bf16、tf16、i8、u8 等。
  • CDFmt：CC、DD 矩阵中元素的精度，可以是 f16、f32、f64、bf16、tf16、i8、u8 等，如果 CDFmt 和ABFmt 一样，则省略 CDFmt。
  • M、N、K：对应上面表述的 A、B、C、D 矩阵的行列数：
    • mA[M][K]: 输入矩阵A；
    • mB[K][N]: 输入矩阵B；
    • mC[M][N]: 累加输入矩阵C；
    • mD[M][N]: 累加结果矩阵D。
  • a：向量寄存器（VGPRs），用于存储输入矩阵 A；
  • b：向量寄存器（VGPRs），用于存储输入矩阵 B；
  • c：向量寄存器（VGPRs），用于存储累加输入矩阵 C；
  • d：向量寄存器（VGPRs），用于存储累加结果矩阵 D。
info

虽然 AMD 也有类似的文档，但其指令格式和 DCU 上还是有很多差异（AMD 用的是 MFMA形式，如 v_mfma_f32_16x16x4_f32(AMD) vs v_mmac_16x16x4_f32(DCU) ），所以本教程适合中国宝宝。

note

我们可以看到，相对于AMD 的指令，DCU 上的指令似乎参数要少。

DCU 上的 Matrix Core 编程（part1）

· 10 min read
genshen
genshen
student in CS

预备知识：Matrix Core 简介​

英伟达GPU上，Volta架构首次引入 tensor core，可以在一个周期内计算完两个4*4矩阵的乘法，其最广泛的用法就是矩阵乘，因此tensor core硬件受到深度学习研究者的广泛欢迎。

AMD 的GPU，也自 CDNA架构开始引入 Matrix Core，来对标 NVIDIA Tensor Core。 由于功能类似，所以有时候，我们也将 Matrix Core 称为 tensor core，但这两者在编程接口上区别还是比较大的。 DCU 架构的GPU，由于部分采用ROCm生态，其上的 Matrix Core 编程接口和 AMD GPU相近，因此很多 AMD GPU的编程资料可供参考。

Fortran77 编译器对 Logical 的处理

· 5 min read
genshen
genshen
student in CS

在Fortran 77 中，可以使用 Logical 类型的变量来表示 .true. 或者 .false.。 实际上，是采用4字节（64位）来存储 Logical。 对于Intel ifort 编译器和 GNU 下的 gfortran 编译器，.true. 的表示会有所不同：

编译器.true..false.
Intel ifort0xFFFFFFFF（全1）0x00000000
GNU gfortran0x000000010x00000000

参考：https://stackoverflow.com/a/61597485