CUDA 版本矩阵乘

说明:html

1.转载请联系本人git

2.代码在最后

问题描述

矩阵乘法 C = aAB + bC
其中a,b为常数,A,B,C为矩阵程序员

实验要求

  1. 根据内存大小测不一样规模矩阵的处理速度(GFLOPS/s),并给出计算公式。
  2. 请计算系统的理论峰值,若是没有达到理论峰值,尝试给出缘由。

方法

CUDA矩阵的优化有多个思路,在本次试验中我使用了shared memory进行访问速度的提高,尝试减小if-else语句的出现,避免串行化,同时作了精度优化以下降错误率(结果不怎么好)。
同时,参考Nvidia给的Samples中0_simple里的matrixMulCUBLAS相关代码,思考提高空间。
github

实验

结果及分析

1.假设矩阵维度为n

处理速度公式=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;
存在着空闲的线程;
以及操做系统的线程调度,和服务器自己的不稳定性等等。服务器

2.优化过程

2.1尝试shared memory

Shared memory的做用在于下降对于全局数据的访问,充分利用Cuda中线程能够有独立的内存空间及寄存器,以及block中线程之间能够通讯的特色
在shared memory大小定义中,Width要保证不能大于XY对应dim的最小值,另外在测试的时候发现,若是width_size大于32,那么获得的结果是全错(不管XY的dim有多大)暂时不清楚为何。
函数

2.2尝试减小if-else语句

在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(),在分配的时候动态设定边界大小,可是参考调用以后其优化的效果不是很明显,没有原做者所说的三倍性能提高,可能和本人的相关知识掌握不足有关。

2.3尝试采用for循环展开

在sgmm函数的for循环以前,使用 #pragma unroll ,GFLOPS/s提高了10个左右的点,效果比较显著。
另外参考课件,有考虑过用Parallel Reduction中的连续访问的方法,可是运行以后程序报错或者错误率很高,暂时没有找到解决办法。

优化后的巅峰状态

2.4 运行CUBLAS对比

CUBLAS 是Nvidia程序员专门优化过的函数,性能表现极好,因为代码不开源,暂时不了解应该如何调整代码。
下面的一次测试显示,在维度为680的矩阵状况下,其performance = 1272 GFlops/s , time = 0.154 msec,较我本身的代码好了三倍有余。

结论

  1. shared memory的正确使用可以很是显著地提高矩阵乘法的性能
  2. if-else语句产生的divergence问题是很是值得关注的,若是不尽可能减小分支语句的使用,并行性能将不会有很好的体现。
  3. CUDA编程的调试难度不亚于OpenMP,以及优化代码须要细心,按部就班。另外也要谨记Amdahl优化定律,抓重要的代码进行性能优化。
  4. CUBLAS的代码表示,可以很是接近峰值是一件很是复杂且辛苦的事情,须要慢慢测量和分析。

参考

《在CUDA中实现任意尺寸的矩阵乘法》

《矩阵乘法—CUDA优化记录》

《利用cuda的cublas库实现任意矩阵的乘法》

代码地址

我的GitHub:Icarusintheworld

相关文章
相关标签/搜索