问题描述
我有一个 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。我注意到与 clblas 或 miopengemm 等其他 Gemm 库类似的行为 - 对于 4096x16 的特定大小,我的性能显着下降,将 N 更改为 1 可多次提高性能。
工作负载分为 256 个线程的工作组。每个工作组处理 128x16 和 128x16 矩阵图块(每个线程 8x8 块)。
我尝试将矩阵平铺改为 96x96 的 6x6 块,而不是 128x128 的 8x8 - 结果相同。
我使用 ROCm 3.7 OpenCL、Clover OpenCL 甚至 Windows OpenCL 驱动程序测试了相同的代码 - 行为相同。
具有相同数量的 GPU 内核(线程)和相同内存类型/大小的 nvidia gtx 960 不存在此类问题。
我怀疑这与缓存/冲突有关,但我不明白它是如何发生的。因此我不知道如何解决它。
解决方法
最后我发现 clBlas 库(最初为 AMD 开发)处理 lda % 1024==0
、ldb % 1024==0
的特殊情况可能是由于缓存
我发现更好的方法是按 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 倍。