CUDA性能调优（一）--合并访问&循环展开

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计数器更新

索引为常量-编译器可以使用加载指令的寻址模式对应的偏移量，这样可以消除地址运算指令。

因此，这个很长的表达式的执行速度几乎可以接近性能的峰值。