加入收藏 | 设为首页 | 会员中心 | 我要投稿 李大同 (https://www.lidatong.com.cn/)- 科技、建站、经验、云计算、5G、大数据,站长网!
当前位置: 首页 > 百科 > 正文

OpenCL代码在MBP上的运行速度比在NVIDIA GTX 480上快

发布时间:2020-12-16 09:34:57 所属栏目:百科 来源:网络整理
导读:我遇到了一个奇怪的问题.到目前为止,我在OpenCL中实现了一些线性代数,只有矩阵乘法,并且已经在我的笔记本电脑上进行了测试.代码非常简单: __kernel void matrix_mult(__global float* a,__global float* b,__global float* c,const int N) { int row = get_
我遇到了一个奇怪的问题.到目前为止,我在OpenCL中实现了一些线性代数,只有矩阵乘法,并且已经在我的笔记本电脑上进行了测试.代码非常简单:

__kernel void matrix_mult(__global float* a,__global float* b,__global float* c,const int N) 
{
  int row = get_global_id(1);
  int col = get_global_id(0);
  float sum = 0.0f;
  for (int i = 0; i < N; i++) {
    sum += a[row*N+i] * b[i*N+col];
  }
  c[row*N+col] = sum;
}

我通过运行代码100次测试硬件,如下所示:

clock_t begin=clock(); 

  const unsigned int repeats = 100;
  for(int  i = 0; i != repeats; i++){
    runCL(a,b,results,N,N*N);
  }

  clock_t end=clock();

在我的MBP上,matrix_multiplications大约需要1.2毫秒,大小为512 * 512,而相同的代码在GTX 480 Linux机器上运行时大约需要3毫秒.这让我感到困扰,我不希望昂贵的GTX卡比笔记本电脑快一点.

据我所知,我的代码是“错误的”我以某种错误的方式计时.

我尝试在OpenCL规范中使用基于事件的计时系统,这给了一些更真实的结果.

cl_event event = {0}; 
err = clEnqueueNDRangeKernel(cmd_queue,kernel[0],2,NULL,global_work_size,&event);
assert(err == CL_SUCCESS);


cl_int err =  clWaitForEvents (1,&event);
cl_ulong start,end; 
clGetEventProfilingInfo(event,CL_PROFILING_COMMAND_END,sizeof(cl_ulong),&end,NULL); 
clGetEventProfilingInfo(event,CL_PROFILING_COMMAND_START,&start,NULL); 
double executionTimeInMilliseconds = (end - start) * 1.0e-6f;
std::cout << "execution time in milis : " << executionTimeInMilliseconds << std::endl;

现在GT330M将在46ms内完成操作,GTX480将在2.5ms内完成.这就产生了另一个非常有趣的问题,在启用PROFILING时,GT 330M的速度变慢约30倍,这种情况很有意义,但GTX480保持了相同的性能.谁能解释为什么会这样?

解决方法

在解决原始问题时,你在这里看到的是,使用这个天真的代码,GTX480的更好的规格实际上伤害了你.

代码样本,矩阵乘法的第一遍,完全由内存带宽支配;每个线程正在访问B的不同元素,由于步幅,它不能被利用.

GTX480的内存总线比GT330M(128位,800 MHz)大3倍(384位)和2倍(1840 MHz).名义上,它提供了177.4GB / s和25.6 GB / s的峰值带宽优势,并且由于这是内存带宽占主导地位,您可能会认为这会赢.但是,由于非合并读取和更宽的内存总线,b阵列访问仅使用该384位存储器访问的32位,而在330M情况下,每128位访问中只有32位.因此,b访问的有效内存带宽为14.8GB / s和6.4GB / s;所以现在总内存带宽只有2倍差异,而不是7左右,而更快的卡的优势就在于浪费;此外,内存带宽必须除以10倍的内核,因此每个内核获得访问权限和进行计算的延迟时间更长.我怀疑如果你使用更大的矩阵大小,你可以隐藏更多的延迟,并接近最佳可能的2倍加速,而不是你看到的2.5倍减速.

这里的最终解决方案是使用更加内存友好的矩阵乘法算法作为基准.

你看到的分析结果,我不知道.也许330M对分析没有那么好的硬件支持,所以事情必须用软件实现?由于GTX数字大致相同,我现在只使用更简单的计时方法,因为你没有使用异步内核或传输,应该没问题.

(编辑:李大同)

【声明】本站内容均来自网络,其相关言论仅代表作者个人观点,不代表本站立场。若无意侵犯到您的权利,请及时与联系站长删除相关内容!

    推荐文章
      热点阅读