OpenCL 矩阵乘法应该更快？

我正在尝试学习如何使 GPU 优化 OpenCL 内核，我以使用本地内存中的方形图块进行矩阵乘法为例。然而在最好的情况下，我只得到了约 10 倍的加速（约 50 Gflops）与 numpy.dot() 相比（5 Gflops，它使用的是 BLAS）。

我发现研究在哪里他们的加速比>200x（>1000 Gflops）. ftp://ftp.u-aizu.ac.jp/u-aizu/doc/Tech-Report/2012/2012-002.pdf ftp://ftp.u-aizu.ac.jp/u-aizu/doc/Tech-Report/2012/2012-002.pdf我不知道我做错了什么，或者只是因为我的 GPU ( nvidia GTX 275 )。或者如果是因为一些 pyOpenCl 开销。但我还测量了将结果从 GPU 复制到 RAM 需要多长时间，它只是矩阵乘法时间的 10% 左右。

#define BLOCK_SIZE 22 
__kernel void matrixMul(
      __global float* Cij, 
      __global float* Aik, 
      __global float* Bkj, 
      __const int ni, 
      __const int nj,
      __const int nk
){
//   WARRNING : interchange of  i  and  j  dimension  lower the performance >2x on my nV GT275 GPU    
int gj = get_global_id(0);    int gi = get_global_id(1); 
int bj = get_group_id(0);     int bi = get_group_id(1);  // Block index
int tj = get_local_id(0);     int ti = get_local_id(1);  // Thread index
int oj = bi*BLOCK_SIZE;       int oi = bj*BLOCK_SIZE; 
float Csub =0; 
__local float As   [BLOCK_SIZE][BLOCK_SIZE];
__local float Bs   [BLOCK_SIZE][BLOCK_SIZE];
for (int ok = 0; ok < nk; ok += BLOCK_SIZE )   {
    As[ti][tj] = Aik[ nk*(gi   ) + tj + ok ];   // A[i][k]
    Bs[ti][tj] = Bkj[ nj*(ti+ok) + gj ];        // B[k][j]
    barrier(CLK_LOCAL_MEM_FENCE);
    for (int k = 0; k < BLOCK_SIZE; ++k) Csub += As[ti][k] * Bs[k][tj];
    barrier(CLK_LOCAL_MEM_FENCE);
}
Cij[ nj * ( gi ) + gj ] = Csub;

}

注意 - 奇怪的 BLOCK_SIZE=22 是最大 BLOCK_SIZE，它适合最大 work_group_size，在我的 GPU 上为 512。在此代码中必须保持条件 BLOCK_SIZE^2

我还尝试了简单的matrixMul（不使用本地内存），但它甚至比numpy.dot()慢10倍。 我在这里复制了代码http://gpgpu-computing4.blogspot.cz/2009/10/matrix-multiplication-3-opencl.html http://gpgpu-computing4.blogspot.cz/2009/10/matrix-multiplication-3-opencl.html他们说即使是简单版本（没有本地内存）的运行速度也应该比 CPU 快 200 倍？我不明白这一点。

在我的例子中，性能的依赖性是：

N =  220 numpy 3.680 [Gflops] GPU 16.428 [Gflops] speedUp 4.464 
N =  330 numpy 4.752 [Gflops] GPU 29.487 [Gflops] speedUp 6.205 
N =  440 numpy 4.914 [Gflops] GPU 37.096 [Gflops] speedUp 7.548 
N =  550 numpy 3.849 [Gflops] GPU 47.019 [Gflops] speedUp 12.217 
N =  660 numpy 5.251 [Gflops] GPU 49.999 [Gflops] speedUp 9.522 
N =  770 numpy 4.565 [Gflops] GPU 48.567 [Gflops] speedUp 10.638 
N =  880 numpy 5.452 [Gflops] GPU 44.444 [Gflops] speedUp 8.152 
N =  990 numpy 4.976 [Gflops] GPU 42.187 [Gflops] speedUp 8.478 
N = 1100 numpy 5.324 [Gflops] GPU 83.187 [Gflops] speedUp 15.625 
N = 1210 numpy 5.401 [Gflops] GPU 57.147 [Gflops] speedUp 10.581 
N = 1320 numpy 5.450 [Gflops] GPU 48.936 [Gflops] speedUp 8.979  

注意 - “Gflops”数字是按 N^3/次获得的，它确实包括将结果从 GPU 复制到主内存所需的时间，但这个时间仅占总时间的百分之几，特别是对于 N>1000

也许更形象的是以秒为单位的时间：

N =  220 numpy 0.003 [s] GPU 0.001 [s] load 0.001 [s] speedUp 5.000 
N =  330 numpy 0.008 [s] GPU 0.001 [s] load 0.001 [s] speedUp 7.683 
N =  440 numpy 0.017 [s] GPU 0.002 [s] load 0.001 [s] speedUp 7.565 
N =  550 numpy 0.043 [s] GPU 0.004 [s] load 0.001 [s] speedUp 11.957 
N =  660 numpy 0.055 [s] GPU 0.006 [s] load 0.002 [s] speedUp 9.298 
N =  770 numpy 0.100 [s] GPU 0.009 [s] load 0.003 [s] speedUp 10.638 
N =  880 numpy 0.125 [s] GPU 0.010 [s] load 0.000 [s] speedUp 12.097 
N =  990 numpy 0.195 [s] GPU 0.015 [s] load 0.000 [s] speedUp 12.581 
N = 1100 numpy 0.250 [s] GPU 0.031 [s] load 0.000 [s] speedUp 8.065 
N = 1210 numpy 0.328 [s] GPU 0.031 [s] load 0.000 [s] speedUp 10.581 
N = 1320 numpy 0.422 [s] GPU 0.047 [s] load 0.000 [s] speedUp 8.979

我在想也许可以使用以下方法来提高速度 async_work_group_copy 甚至 read_imageui 将块复制到本地内存。但我不明白为什么当我使用与那些说他们有 200 倍加速的人基本相同的代码时我有这么大的差异？？？

甚至不看你的代码，让我对你的基准进行一些评论。让我们忽略 numpy，并比较 Intel CPU 与 Nvidia 和 AMD GPU 的最大 SP FLOPs/s 和 DP FLOPs/s。

4 GHz 的 Intel 2600K 可以执行 4 GHz * (8 AVX) * (2 ILP) * ( 4 核) = 256 SP GFLOPs/s。对于 DP，它是一半：128 DP GFLOPs/s。几周后推出的 Haswell 将使这两个数字翻倍。英特尔 MKL 库在 GEMM 中的效率高于 80%。我自己的 GEMM 代码在我的 i7-2700 上获得了 70%，所以你用 numpy 引用的 5 GFlops/s 很小，比较不公平。

我不知道 GTX 275 的能力如何，但我猜它的速度远远超过 50 GFLOPs/s。

您引用的文章对 AMD 7970 进行了比较。它们获得 848（90% 效率）DP GFlops/s 和 2646（70% 效率）SP GFlops/s。这接近 CPU 性能的 10 倍，而不是 200 倍！

编辑： 您对 FLOP 的计算是错误的，它应该是 2.0*n^3。这仍然是近似值，但渐近正确。让我解释。

考虑 3D 点积。是x1*x2+y1*y2+z1*z2。这是 3 次乘法和 2 次加法。因此，N 维点积是 n 次乘法和 (n-1) 次加法。矩阵乘积相当于 nxn 点积，即 n*n*n 次乘法和 n*n*(n-1) 次加法。这大约是 2.0*n^3 FLOPS。因此，您应该将所有 Gflops/s 数字加倍。

编辑： 您可能需要考虑内核时间。自从我使用 OpenCL 以来已经有一段时间了，但是使用 C++ 绑定我做了这样的事情

queue = cl::CommandQueue(context, devices[device], CL_QUEUE_PROFILING_ENABLE|CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err);
//other code...run kernel

time_end = clevent.getProfilingInfo<CL_PROFILING_COMMAND_END>();  
time_start = clevent.getProfilingInfo<CL_PROFILING_COMMAND_START>();