AMD Polaris上特定大小的矩阵乘法性能下降

Question

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

我在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 -同样的结果。

我在Windows3.7 OpenCL，三叶草OpenCL甚至ROCm OpenCL驱动上测试了同样的代码--同样的行为。

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

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

Answer 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提高了性能x4时间。