CUDA性能调优(一)--合并访问&循环展开
2017-03-25 00:19
162 查看
1 合并访问
当同一个warp中的所有线程都执行同一条指令访问全局存储器中连续的单元时,就获得最有利的访问模式。硬件检测到同一个warp中的这些线程访问全局存储器中连续的存储单元,并将这些单元结合成一个合并的访问。
合并访问可以提高DRAM的带宽利用率,使DRAM在传输数据时的速度接近全局存储器带宽的峰值。
这里需要补充一点知识:二维三维数组的线性映射--多维线程,维度映射到线性顺序。
二维线程按照线性顺序排列如下图所示:
![](https://img-blog.csdn.net/20170325104349549)
也就是说矩阵在内存空间中已经被等价线性化为等价的一维数组,以行优先的方式逐行顺序存放。
示例如下:
(1) 在未[b]使用shared memory的GPU矩阵相乘中的两种访问模式[/b]
![](https://img-blog.csdn.net/20170325095759065)
接合访问模式
在第一次迭代中,访问0号元素,这些0号元素在global memory中相邻,硬件接合这些访问:
![](https://img-blog.csdn.net/20170919184211692)
非接合访问模式
![](https://img-blog.csdn.net/20170919184236575)
在这个例子中得出如下结论:
对于全局存储器的访问,kernel循环遍历一行时,要比遍历一列的效率低得多。
(2) 在使用shared memory的矩阵相乘中,使用合并方式加载数据
每个线程负责加载一个Md及Nd元素。
m识别左侧瓦片的位置。
Md瓦片的一行由TILE_WIDTH个线程加载,threadIdx.x变化。
2 指令混合
在当前设备中,每个SM的指令处理带宽有限。每个指令占用指令处理带宽,包括浮点计算指令、加载指令和分支指令。消除重复指令可减少指令处理带宽的压力,提升kernel函数的整体执行性能。
以下面两行代码为例进行介绍:
(1)循环引入额外指令更新计数器k 1次
(2)每次迭代结尾位置执行条件跳转 1次
(3)使用k计算Mds,Nds索引引入了地址运算指令 2次
(4)浮点乘加计算指令 2次
浮点乘加计算指令仅占了1/3的指令。由于指令处理带宽有限,因此这种指令混合将能取得的性能限制在带宽峰值的1/3以内。
为了改善这用指令混合,采取循环展开(Unroll loop)的方法。
将上面的代码修改如下:
代码分析:
Long-multiply-add操作
消除分支指令以及loop计数器更新
索引为常量-编译器可以使用加载指令的寻址模式对应的偏移量,这样可以消除地址运算指令。
因此,这个很长的表达式的执行速度几乎可以接近性能的峰值。
当同一个warp中的所有线程都执行同一条指令访问全局存储器中连续的单元时,就获得最有利的访问模式。硬件检测到同一个warp中的这些线程访问全局存储器中连续的存储单元,并将这些单元结合成一个合并的访问。
合并访问可以提高DRAM的带宽利用率,使DRAM在传输数据时的速度接近全局存储器带宽的峰值。
这里需要补充一点知识:二维三维数组的线性映射--多维线程,维度映射到线性顺序。
二维线程按照线性顺序排列如下图所示:
也就是说矩阵在内存空间中已经被等价线性化为等价的一维数组,以行优先的方式逐行顺序存放。
示例如下:
(1) 在未[b]使用shared memory的GPU矩阵相乘中的两种访问模式[/b]
接合访问模式
在第一次迭代中,访问0号元素,这些0号元素在global memory中相邻,硬件接合这些访问:
非接合访问模式
在这个例子中得出如下结论:
对于全局存储器的访问,kernel循环遍历一行时,要比遍历一列的效率低得多。
(2) 在使用shared memory的矩阵相乘中,使用合并方式加载数据
//以合并的模式加载 ds_A[ty][tx] = A[Row*n+t*TILE_WIDTH+tx]; ds_B[ty][tx] = B[(t*TILE_WIDTH + ty)*k+Col];
每个线程负责加载一个Md及Nd元素。
m识别左侧瓦片的位置。
Md瓦片的一行由TILE_WIDTH个线程加载,threadIdx.x变化。
2 指令混合
在当前设备中,每个SM的指令处理带宽有限。每个指令占用指令处理带宽,包括浮点计算指令、加载指令和分支指令。消除重复指令可减少指令处理带宽的压力,提升kernel函数的整体执行性能。
以下面两行代码为例进行介绍:
for (int i = 0; i < TILE_WIDTH; ++i) Cvalue += ds_A[ty][i] * ds_B[i][tx]这两行代码包括一下几种指令:
(1)循环引入额外指令更新计数器k 1次
(2)每次迭代结尾位置执行条件跳转 1次
(3)使用k计算Mds,Nds索引引入了地址运算指令 2次
(4)浮点乘加计算指令 2次
浮点乘加计算指令仅占了1/3的指令。由于指令处理带宽有限,因此这种指令混合将能取得的性能限制在带宽峰值的1/3以内。
为了改善这用指令混合,采取循环展开(Unroll loop)的方法。
将上面的代码修改如下:
Cvalue += ds_A[ty][0] * ds_B[0][tx]+ds_A[ty][1] * ds_B[1][tx] +ds_A[ty][2] * ds_B[2][tx]+ds_A[ty][3] * ds_B[3][tx] +ds_A[ty][4] * ds_B[4][tx]+ds_A[ty][5] * ds_B[5][tx] +ds_A[ty][6] * ds_B[6][tx]+ds_A[ty][7] * ds_B[7][tx] +ds_A[ty][8] * ds_B[8][tx]+ds_A[ty][9] * ds_B[9][tx] +ds_A[ty][10] * ds_B[10][tx]+ds_A[ty][11] * ds_B[11][tx] +ds_A[ty][12] * ds_B[12][tx]+ds_A[ty][13] * ds_B[13][tx] +ds_A[ty][14] * ds_B[14][tx]+ds_A[ty][15] * ds_B[15][tx];
代码分析:
Long-multiply-add操作
消除分支指令以及loop计数器更新
索引为常量-编译器可以使用加载指令的寻址模式对应的偏移量,这样可以消除地址运算指令。
因此,这个很长的表达式的执行速度几乎可以接近性能的峰值。
相关文章推荐
- 菜单展开 合并<二>
- CUDA:对齐内存访问、展开循环提高运算性能
- 利用IList<T>同时循环访问列表和数组
- Java API —— IO流(数据操作流 & 内存操作流 & 打印流 & 标准输入输出流 & 随机访问流 & 合并流 & 序列化流 & Properties & NIO)
- 菜单展开 合并<一>
- 修复"打开方式"时会弹出"拒绝访问"的问题
- "SQL Server不存在或访问被拒绝"问题的解决
- asp + Access 常见的数据库访问失败问题 Microsoft JET Database Engine 错误 '80004005' 解决办法
- JTree展开结点设置为'+'"-"形式
- 怎样解决:对路径""的访问被拒绝 这个问题?
- 解决访问优盘(MP3)出现 "本次操作由于受计算机的限制而被取消,请与你的系统管理员联系" 的问题
- js入门·循环与判断/利用函数的简单实例/使用对象/列举对象属性的名称
- 一个网页展开合并的小例子
- 利用Windows消息循环实现消息延迟和消息合并
- msxml3.dll 错误 '80070005' 拒绝访问。
- "访问由于凭据无效被拒绝"的解决办法
- Ntfs 分区的硬盘卷删除了权限重新恢复后"拒绝访问"![解决办法]
- 单循环比赛的"贝格尔"编排法
- 打开文件->选择程序 出现"拒绝访问"
- 在XSL里执行类似for(i=0;i<n;i++)的循环----我的实现