Skip to main content

5 posts tagged with "hpc"

View All Tags

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的编程资料可供参考。

在 CUDA 环境下安装 HIP

· 9 min read
genshen
genshen

HIP (Heterogeneous-Compute Interface for Portability) 是 AMD 开发的一款异构计算的接口工具。 HIP 允许只用写一套代码(hip代码), 就可以将程序同时在 NVIDIA GPU 和 AMD GPU 及 DCU 上编译运行。

HIP is a C++ Runtime API and Kernel Language that allows developers to create portable applications for AMD and NVIDIA GPUs from single source code.

HIP 的 API 和 CUDA 的API十分类似，例如 CUDA 中内存拷贝用cuMemory, 在 hip 中用hipMemcpy，且参数也十分一致。 因此，会 CUDA 的开发者可以很轻松地转移到 hip 上。 并且，hip 还提供了hipfy 工具，将 CUDA 代码转换为 hip 代码。
HIP 在不降低性能的前提下，统一了CUDA API 和AMD GPU 编程API，可谓极大地降低了各个平台的适配与移植工作， 做到了一套代码，在多个异构平台上运行。 可以说, "舍弃 CUDA，进入HIP时代"。

那么，在 NV GPU下，如何安装并使用 hip 呢？