说明:html
1.转载请联系本人git
2.代码在最后矩阵乘法 C = aAB + bC
其中a,b为常数,A,B,C为矩阵程序员
CUDA矩阵的优化有多个思路,在本次试验中我使用了shared memory进行访问速度的提高,尝试减小if-else语句的出现,避免串行化,同时作了精度优化以下降错误率(结果不怎么好)。
同时,参考Nvidia给的Samples中0_simple里的matrixMulCUBLAS相关代码,思考提高空间。
github
处理速度公式=2*n/1000000000/time;
带宽计算公式:
= ( sizeof(int)*dim + sizeof(int)*n + sizeof(float)*n
+ sizeof(float)*dim*2)/1000000000/time;编程
系统理论峰值(即浮点数理论峰值)
集群理论浮点峰值
= CPU主频(GHz)× CPU每时钟周期执行浮点运算次数 × 节点数 × 8(每节点双路四核)
=4.2*4*8=134.4GFLOPS/s
性能优化
峰值带宽: B=F×D/8=2133MHz*64bit/8=17.064GHz
bash
没有达到理论峰值的缘由是:
程序并不仅是在作浮点数运算或只是在访问内存;
sgemm中还存在着if-else语句,使得线程存在着divergence;
因为大小分配的问题存在着Occupancy;
存在着空闲的线程;
以及操做系统的线程调度,和服务器自己的不稳定性等等。服务器
Shared memory的做用在于下降对于全局数据的访问,充分利用Cuda中线程能够有独立的内存空间及寄存器,以及block中线程之间能够通讯的特色
在shared memory大小定义中,Width要保证不能大于XY对应dim的最小值,另外在测试的时候发现,若是width_size大于32,那么获得的结果是全错(不管XY的dim有多大)暂时不清楚为何。
函数
在Sgemm函数中,if-else语句主要用于进行边界判断。
这是由于在分配block大小的时候,矩阵的维度可能不能恰好被32整除。例如dim=500时,不进行边界判断会引发不少问题。
一个有效的解决方案是,利用ceil的取整函数,在for循环中有效限制i的上界。使得对矩阵维度的限制没有那么大。性能
在代码中对grid, block 定义以下
dim3 block(DIM_THREAD_BLOCK_Y, DIM_THREAD_BLOCK_Y);
dim3 grid((size_t)ceil( ((float)N) / ((float)block.x) ),
(size_t)ceil( ((float)N) / ((float)block.y)) );
//取整函数ceil
复制代码
固然,通过反复测试代表,矩阵的维度若能被32整除,其性能表现要比不能整除的要好。
另外在搜索查找的时候看到有一个方式是利用了cudaMallocPitch(),在分配的时候动态设定边界大小,可是参考调用以后其优化的效果不是很明显,没有原做者所说的三倍性能提高,可能和本人的相关知识掌握不足有关。
在sgmm函数的for循环以前,使用 #pragma unroll ,GFLOPS/s提高了10个左右的点,效果比较显著。
另外参考课件,有考虑过用Parallel Reduction中的连续访问的方法,可是运行以后程序报错或者错误率很高,暂时没有找到解决办法。
优化后的巅峰状态
CUBLAS 是Nvidia程序员专门优化过的函数,性能表现极好,因为代码不开源,暂时不了解应该如何调整代码。
下面的一次测试显示,在维度为680的矩阵状况下,其performance = 1272 GFlops/s , time = 0.154 msec,较我本身的代码好了三倍有余。
我的GitHub:Icarusintheworld