OpenCL代码在MBP上的运行速度比在NVIDIA GTX 480上快
我遇到了一个奇怪的问题.到目前为止,我在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数字大致相同,我现在只使用更简单的计时方法,因为你没有使用异步内核或传输,应该没问题. (编辑:李大同) 【声明】本站内容均来自网络,其相关言论仅代表作者个人观点,不代表本站立场。若无意侵犯到您的权利,请及时与联系站长删除相关内容! |
- gxpt之规范自定义xml配置——xsd显威
- 如果sqlite安装在非标准位置,则构建python 2.6 w
- Client found response content type of 'mu
- 快速解决C# android base-64 字符数组的无效长度
- u-boot for tiny210 ver3.0 (by liukun321咕唧咕
- XML文档的四种生成和解析方法详解
- c# – 从不执行的可执行文件中获取版本
- c# – Visual Studio Azure函数发布function.jso
- c# – 我如何模拟这个异步方法?
- 关于Oracle 12c新特性---Rapid Home Provisionin