1 合并访问
当同一个warp中的所有线程都执行同一条指令访问全局存储器中连续的单元时,就获得最有利的访问模式。硬件检测到同一个warp中的这些线程访问全局存储器中连续的存储单元,并将这些单元结合成一个合并的访问。
合并访问可以提高DRAM的带宽利用率,使DRAM在传输数据时的速度接近全局存储器带宽的峰值。
这里需要补充一点知识:二维三维数组的线性映射--多维线程,维度映射到线性顺序。
二维线程按照线性顺序排列如下图所示:
也就是说矩阵在内存空间中已经被等价线性化为等价的一维数组,以行优先的方式逐行顺序存放。
示例如下:
(1) 在未使用shared memory的GPU矩阵相乘中的两种访问模式
接合访问模式
在第一次迭代中,访问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计数器更新
索引为常量-编译器可以使用加载指令的寻址模式对应的偏移量,这样可以消除地址运算指令。
因此,这个很长的表达式的执行速度几乎可以接近性能的峰值。