AMD Polaris 上某些尺寸的矩阵乘法性能下降
Performance drop in matrix multiplication for certain sizes on AMD Polaris
我有一个将 2 个矩阵 (GEMM) 与 M=4096、N=4096 和 K=16 相乘的 OpenCL 代码。 （即矩阵 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 没有这样的问题 type/size。
我怀疑这在某种程度上与 cache/collision 有关，但我不明白它是如何发生的。因此我不知道如何解决它。
最后我发现 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 顺序性能提高了 x4 倍。
我有一个将 2 个矩阵 (GEMM) 与 M=4096、N=4096 和 K=16 相乘的 OpenCL 代码。 （即矩阵 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 没有这样的问题 type/size。
我怀疑这在某种程度上与 cache/collision 有关，但我不明白它是如何发生的。因此我不知道如何解决它。
最后我发现 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 顺序性能提高了 x4 倍。