我有一个 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 块）。
我尝试使用 6x6 块将矩阵平铺更改为 96x96，而不是使用 8x8 的 128x128 - 结果相同。
我使用 ROCm 3.7 OpenCL、Clover OpenCL 甚至 Windows OpenCL 驱动程序测试了相同的代码 - 相同的行为。
具有相同数量的 gpu 内核（线程）和相同内存类型/大小的 nvidia gtx 960 不存在此类问题。
我怀疑这与缓存/冲突有关，但我不明白它是如何发生的。因此，我不知道如何解决它。