AMD Polaris 上某些尺寸的矩阵乘法性能下降答案

【问题标题】：Performance drop in matrix multiplication for certain sizes on AMD PolarisAMD Polaris 上某些尺寸的矩阵乘法性能下降
【发布时间】：2021-09-09 22:12:17
【问题描述】：

我有一个 OpenCL 代码，它将 2 个矩阵 (GEMM) 与 M=4096、N=4096 和 K=16 相乘。（即矩阵 4096 x 16 浮点数）

我在 Polaris 560 16CU GPU 上运行它。

代码：https://github.com/artyom-beilis/oclblas/blob/master/gemm/gemm.cl

我注意到这种大小的性能下降非常奇怪，这种大小的矩阵乘法具有大约 8-10 GFlops 的性能，而如果我将 N 更改为 4095 或 4097，我会得到大约 130-150Gflops 的性能。我注意到与其他 GEMM 库（如 clblas 或 miopengemm）类似的行为 - 对于这种 4096x16 的特定大小，我的性能显着下降，将 N 更改为 1 可将性能提高数倍。

工作负载分为 256 个线程的工作组。每个工作组处理 128x16 和 128x16 矩阵图块（每个线程 8x8 块）。

我尝试将矩阵平铺更改为 96x96 和 6x6 块，而不是 128x128 和 8x8 - 结果相同。

我使用 ROCm 3.7 OpenCL、Clover OpenCL 甚至 Windows OpenCL 驱动程序测试了相同的代码 - 行为相同。

具有相同数量的 gpu 内核（线程）和相同内存类型/大小的 nvidia gtx 960 不存在此类问题。

我怀疑这在某种程度上与缓存/冲突有关，但我不明白它是如何发生的。因此我不知道如何解决它。

【问题讨论】：

我对矩阵乘法的不同主题做了相当多的研究。这篇 stackoverflow 文章可以带您走上正确的道路：stackoverflow.com/questions/31439512/… 看看，让我知道您的想法。希望我们可以为此评估适当的解决方案。我注意到 ATI/AMD 显卡没有像 NVIDEA 那样优化。帧缓冲区问题和锁定线程会导致 gflops 下降。
@DEX7RA 高效 GEMM 的一般概念已经实现 - 更具体的是某种我无法理解的内存访问模式缓存问题
我明白了，但我发现一些高水平的研究表明 Gflops 下降是正常的。如果它们不一致，但有时会发生不规则的情况。这就解释了为什么它们存在于更小甚至更大的矩阵中。
是的，但问题是性能下降了一个数量级

标签： gpu opencl matrix-multiplication amd-gcn

【解决方案1】：

最后我发现 clBlas 库（最初是为 AMD 开发的）处理 lda % 1024==0、ldb % 1024==0 的特殊情况可能是由于缓存

https://github.com/clMathLibraries/clBLAS/blob/master/src/library/blas/specialCases/GemmSpecialCases.cpp#L228

我发现更好的方法是按 z 曲线顺序重新排列块，而不是让多个内核排队。

https://github.com/artyom-beilis/oclblas/blob/master/gemm/gemm.cl#L109

为了处理 M!=N 或 M != 1<<n 的案例，我只是将 M/N 上的工作组数量增加到接近 1<<n 并且没有工作的组在乞求中退出而不增加太多开销。

z-order 将性能提高了 4 倍。

【讨论】：

哇，很酷，你发现了！对不起，我没有早点回复，但没有任何线索，正在研究。但是你更快。很高兴它有效！