OpenCL矩阵乘法应该更快？

Question

我正在尝试学习如何制作GPU优化的OpenCL kernells,我在本地存储器中使用方形图块作为矩阵乘法的例子.然而,与numpy.dot()(5 Gflops,它正在使用BLAS)相比,我获得了最好的情况只有~10倍的加速(~50 Gflops).

我发现他们的研究速度超过200倍(> 1000 Gflops). 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,它适合我的GPU上的最大work_group_size,即512.在此代码中必须保持条件BLOCK_SIZE ^ 2 <max work_group_size.22 = INT(SQRT(512)).我也尝试过BLOCK_SIZE = 16或8但是它的速度比较慢.

我也尝试过简单的matrixMul(不使用本地内存),但它甚至比numpy.dot()慢10倍.我在这里复制了代码 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

也许更多的画面是时间在secons:

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倍加速的人时,我有这么大的差异？

Answer 1

如果没有查看您的代码,请让我对您的基准做一些评论.让我们忽略numpy并比较Intel CPU与Nvidia和AMD GPU的最大SP FLOP/s和DP FLOP/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次乘法和两次加法.因此,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>();