AMD OpenCL大学教程(9)

    本节主要讲述GPU的memory架构。优化基于GPU device的kernel程序时，我们需要了解很多GPU的memory知识，比如内存合并，bank conflit（冲突）等等，这样才能针对具体算法做一些优化工作。

1、GPU总线寻址介绍

   假定X是一个指向整数（32位整数）数组的指针，数组的首地址为0x00001232。一个线程要访问元素X[0],

   int tmp = X[0];

  

    假定memory总线宽度为256位(HD5870就是如此,即为32字节），因为基于字节地址的总线要访问memory，必须和总线宽度对齐，也就是说按必须32字节对齐来访问memory，比如访问0x00000000,0x00000020,0x00000040,…等，所以我们要得到地址0x00001232中的数据，比如访问地址0x00001220,这时，它会同时得到0x00001220到 0x0000123F 的所有数据。因为我们只是取的一个32位整数，所以有用的数据是4个字节，其它28的字节的数据都被浪费了，白白消耗了带宽。

   

2、合并内存访问

    为了利用总线带宽，GPU通常把多个线程的内存访问尽量合并到较少的内存请求命令中去。

    假定下面的OpenCL kernel代码：int tmp = X[get_global_id(0)];

数组X的首地址和前面例子一样，也是0x00001232，则前16个线程将访问地址：0x00001232 到 0x00001272。假设每个memory访问请求都单独发送的话，则有16个request，有用的数据只有64字节，浪费掉了448字节（16*28）。

    假定多个线程访问32个字节以内的地址，它们的访问可以通过一个memory request完成，这样可以大大提高带宽利用率，在专业术语描述中这样的合并访问称作coalescing。

   例如上面16个线程访问地址0x00001232 到 0x00001272，我们只需要3次memory requst。

   在HD5870显卡中，一个wave中16个连续线程的内存访问会被合并，称作quarter-wavefront，是重要的硬件调度单位。

   下面的图是HD5870中，使用memory访问合并以及没有使用合并的bandwidth比较：

   下图是GTX285中的比较：

3、Global memory的bank以及channel访问冲突

   我们知道内存由bank，channel组成，bank是实际存储数据的单元，一个mc可以连接多个channel，形成单mc，多channel的连接方式。在物理上，不同bank的数据可以同时访问，相同的bank的数据则必须串行访问，channel也是同样的道理。但由于合并访问的缘故，对于global memory来说，bank conflit影响要小很多，除非是非合访问，不同线程访问同一个bank。理想情况下，我们应该做到不同的workgroup访问的不同的bank，同一个group内，最好用合并操作。

   

4、local memory的bank conflit

   bank访问冲突对local memory操作有更大的影响（相比于global memory），连续的local memory访问地址，应该映射到不同的bank上，

     在AMD显卡中，一个产生bank访问冲突wavefront将会等待所有的local memory访问完成，硬件不能通过切换到另一个wave来隐藏local memory访问时延。所以对local memory访问的优化就很重要。HD5870显卡中，每个CU（simd）有32bank，每个bank 1k，按4字节对齐访问。如果没有bank conflit，每个bank能够没有延时的返回一个数据，下面的图就是这种情况。

   如果多个memory访问对应到一个bank上，则conflits的数量决定时延的大小。下面的访问方式将会有3倍的时延。

  但是，如果所有访问都映射到一个bank上，则系统会广播数据访问，不会产生额外时延。