CUDA 版本矩阵乘

说明:

1.转载请联系本人

2.代码在最后

问题描述

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

实验要求

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

方法

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

实验

结果及分析
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

没有达到理论峰值的原因是:
程序并不只是在做浮点数运算或只是在访问内存;
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

  • 1
    点赞
  • 1
    收藏
    觉得还不错? 一键收藏
  • 0
    评论

“相关推荐”对你有帮助么?

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值