OpenCL AMD Cayman架构GPU|极客笔记

AMD Radeon HD Graphics从之前的R700架构开始就能支持OpenCL 1.0了。R700到Evergreen架构全都采用了VLIW5的执行引擎。而到了Radeon HD 6900系列，则采用了TeraScale3架构的VLIW4的执行引擎。而现在从GCN架构起，都采用了标量处理单元。什么是VLIW（Very Long Instruction Words），什么又是标量处理单元？这两个概念将会在8.1.1节和8.1.2节进行解释。
由于现在VLIW的执行架构已经逐渐被AMD所淘汰，而NVIDIA支持CUDA的统一渲染架构的GPU从一开始就用了标量执行方式。不过就VLIW方式的独特性而言，这里还是可以给大家分享一下。而下面就以Radeon HD Graphics 6900系的VLIW4引擎作为例子给大家介绍一下它的整体结构以及执行方式。

AMD Cayman架构GPU

Cayman的整体架构与执行模型

下图为AMD Radeon HD Graphics 6900系GPU（Cayman架构）的整体架构图。

Cayman架构图

上图中，最顶上的橙色横条是命令处理器（Command Processor），这个用于处理从主机端通过clEnqueue系API发送给GPU设备端的命令。然后，我们看到6900系GPU具有双图形引擎（通用计算引擎也包含在其中），每个图形引擎对应一整块SIMD引擎。图形引擎到SIMD引擎之间还包含了超大规模线程调度处理器，这个用于确定将当前GPU指令分派给哪些SIMD处理单元进行执行。在它下面还含有一个指令Cache和一个常量Cache。GPU在执行内核程序时往往是直接从指令Cache进行取指令，然后再派发到各个SIMD处理单元的。
我们下面简单看一下比较关键的SIMD引擎。SIMD引擎部分由若干SIMD处理单元组成。对于Cayman架构的GPU而言，最多可拥有24个SIMD处理单元。图8-1中，我们就拿右边的SIMD引擎这个区域图为例来说明一下。最右边的是局部数据共享存储器（Local Data Share，LDS），每个SIMD处理单元专用一块。然后，中间四个红色块（每个红色块里包含4个黑色块）就是16个流处理器（Stream Processor，SP），再左边黄色的4块是纹理单元，也是每个SIMD处理单元独用的，然后最左边的紫色块是L1纹理Cache。这一整条块共同组成了一个SIMD处理单元。而我们数一下，左边和右边的SIMD引擎各有12个SIMD处理单元，加起来正好是24个。在Cayman架构的GPU中，一个SIMD处理单元就对应了OpenCL中的计算单元（Compute Unit，CU）。而一个SP则对应了OpenCL中的处理元素（Processing Element，PE）。下面，我们看一下每个SP的内部结构，如下图所示：

SP的内部结构

这是整个流处理器的结构。里面包含了4个流处理单元（Stream Processing Unit）和一个分支处理单元。对于普通的整数或单精度浮点计算来说，可以同时使用这4个流处理单元进行并行操作。当执行三角函数等超越函数的计算时，需要占用3个流处理单元执行计算。所以，对于Cayman架构的GPU而言，使用int4、float4等向量数据进行计算是非常合适的，这能最大化地利用此GPU的VLIW4计算特性。
VILW的执行方式与SIMD的执行方式略有区别。SIMD是单指令多数据。也就是说，对于一个包含了4条f loat数据类型的SIMD寄存器而言，做一次f loat4的加法计算是对该寄存器所有通道（lane）执行一次加法计算，如下图所示：

SIMD示例图

上图中，有两个SIMD寄存器vector0和vector1做一个向量操作，假设是vector0=vector0+vector1。每个SIMD寄存器含有4个分量（component），每个分量作为单独的一个数据元素（如果是f loat4的一个向量，那么每个分量就是一个f loat数据），也被称为标量。在很多书上会使用lane（通道）来描述一个分量。在本例中，分别使用x、y、z和w来分别表示这4个分量，它们彼此相互独立。当执行vector0=vector0+vector1的向量加法时，相当于对vector0和vector1所有4个分量同时做加法计算，即vector0.x +=vector1.x；vector0.y += vector1.y；vector0.z += vector1.z；vector0.w += vector1.w；。这个操作同时进行。这就是所谓的SIMD操作模式，即一条指令对多组数据同时操作。从宏观上讲， Cayman架构GPU的执行方式也是SIMD操作方式，不过对于每个流处理器而言，在指令的执行上又是通过VLIW4执行的。
VLIW4又是怎样的情况呢？我们看下图：

示意图

上图中，整个矩形是一条长指令，它被分为4个slot（指令槽），每个槽是一个简单的指令操作。我们看到，槽0是将reg0与reg1寄存器做加法运算；槽1是将reg2和reg3寄存器做减法计算；槽2是将reg4和reg5寄存器做按位与计算；槽3是将reg6与reg7寄存器做按位或计算。由于每个槽所使用的寄存器都完全相互独立，而且每个槽所使用的算术逻辑操作都十分简单，从而只需要一个ALU算术逻辑单元就能执行了，因此可以把这4个简单的操作拼成一个大的指令，同时执行这4个简单的操作。这就是VLIW4的执行模式。在计算机术语中这也称为MIMD执行模式（多指令多数据）。而在Cayman架构的GPU中，一个SP中甚至可以有两对存在相互依赖的长指令。
对比SIMD和VLIW来说，其执行的区别在于VLIW单元中每个slot可以执行不同指令（在硬件上，指令类型通常有限制），而SIMD单元中每个slot必须执行共同的指令。VLIW单元一定能够执行SIMD操作，而反之不行。
以上介绍了Cayman架构GPU的大致结构以及执行单元特征。下面就来谈谈它是如何执行程序的。Cayman架构的GPU是以wavefront的方式执行的。wavefront可理解为在GPU上执行的一个线程，它可以对连续的64个工作项进行操作，但是需要分4个周期才能完成。也就是说，一个周期只完成1/4个wavefront，即16个工作项。这也是与Cayman架构有关，一个SIMD单元中含有16个SP，每个SP对应执行一个工作项。不过，一个SIMD其实并不是只对应一个工作组，而是可对应多达8个工作组，一共32条wavefront（每个工作组含有4条wavefront）可用于调度。这其实也是为了尽可能地去掩盖访问外部存储器所带来的巨大延迟。但是在执行时一次只调度一条wavefront进行执行，4个周期完成后再去调度其他的wavefront。因此，一条wavefront在Cayman架构中也被称为一个分支粒度（branch granularity）。所以，从应用程序员的角度把它看作为一个原子的线程也没什么问题。
前面已经提到了，Cayman架构的GPU是以VLIW4模式执行的，因此对于一个SIMD处理单元而言，执行一条指令时在最好的情况下是4个流处理单元能同时执行。而一整条wavefront执行完时，一共处理完64个工作项，而最多可完成256个像32位整数加法、单精度浮点加法等这样简单的指令操作。
正是因为我们要充分利用VLIW4的执行特性，所以我们在为Radeon HD 6900系GPU编写OpenCL的内核程序时，应当尽量采用向量数据操作方式。比如int4、f loat4这种数据类型。因为，VLIW4的指令组合其实是通过OpenCL编译器来做的。而对于通用计算程序而言，很难在某个上下文中找到前后相互独立的简单操作可以拼成一条长指令。而向量操作是很天然的组成VLIW4的方式。比如：假设定义了f loat4 a,b；做a+=b；时，其实编译器会把这条向量加法拆分成a.x+=b.x；a.y+=b.y；a.z+=b.z；a.w+=b.w；，然后将这4个操作分别作为一条长指令的各个指令槽，从而能获得4个加法操作并行执行的能力。如果编译器在当前指令上下文中只能操作单个数据，那么此时，SP只能利用它其中一个SPU进行执行，从而性能将会降低到峰值的1/4。
以上就是对Cayman架构GPGPU的简单描述。这里再要提醒各位读者的是，该GPU的最小并行粒度就是64个工作项，即一条完整的wavefront。因此，我们在对该系列的GPU分配工作组时要注意，工作组大小尽量是64的倍数，这样能充分发挥wavefront的计算优势。同时，一条wavefront也作为该GPU的分支粒度。也就是说，如果在一条wavefront内的工作项发生了不同分支执行，比如：

if(index 〈 32)
    func1();
else
    func2();

那么在执行前32个工作项时，后32个工作项也会做func1（）分支，只不过func1（）调用之后不会产生影响（基于GPU的谓词预测（Predicate）执行机制，我们也可以把它当做CPU上的带条件语句执行，如x86上的CMOV），但后32个工作项肯定会等前32个工作项都完成func1（）的执行之后，这64个工作项再会一起往下走。而走到else这个分支也是如此，前32个工作项也会执行func2（），但不会做实际对func2（）调用的效果，不过肯定需要等后32个工作项都执行完func2（）之后才一起往下走。

Cayman架构GPU的LDS结构与优化

下面来谈谈关于Cayman架构GPU的局部存储器访问的优化。对于这一系列的GPU，每个SIMD单元含有一个32KB的局部存储器（Local Data Share，LDS），对应于OpenCL中的局部存储器（Local Memory），我们从图8-1中也能看到。所以，在OpenCL中局部存储器只能被当前工作组中的所有工作项共享，而不能跨工作组进行访问。该32KB的LDS含有32个存储体（bank）。每个存储体具有4字节的宽度以及256字节的深度，这样一个存储体的容量正好是1KB。下面给出一个简单的LDS的结构图，如下图所示。

LDS的结构图

在一个周期内，LDS可以为每个存储体处理一个请求。也就是说，如果当前每个工作项访问各自索引的局部存储器，然后考虑到Radeon HD Graphics 6900系每个周期给16个工作项发射一条指令，那么这16个工作项对局部存储器的访问在此周期内将被并行处理。由于LDS一共有32个存储体，因此如果想充分利用LDS的带宽峰值，每个工作项可以访问8字节的LDS数据，即使用两个存储体。那么16个工作项正好用满了32个存储体。
而如果16个工作项中有任意两个工作项不小心访问了同一个LDS存储体，那么就会发生存储体冲突。此时LDS需要花费更多额外的周期去检查并解决冲突，不管怎么说访问同一个存储体的两个请求会被先后按顺序处理，而不能被并行访问。所以，最糟糕的情况就是16个工作项都同时访问了同一个存储体，那么这16次对LDS的访问都将被串行执行，从而造成LDS访问带宽的最低性能。而这里有一个例外，如果所有16个工作项都访问了同一个LDS的地址（注意，同一个地址与同一个存储体是两个不同的概念），那么所有请求都会被广播，而不会导致存储体冲突。下面列出一个具体的代码示例来讲解LDS存储体冲不冲突的情况。

     //这里使用32KB的LDS空间,正好用满一个CU（即SIMD单元）的整个LDS
     //32个存储体,每个存储体256行的深度,4字节的宽度
     _local int ldsMemory[32 ＊ 256];
     int index = get_global_id(0);
     //每个工作项访问各自4字节的LDS空间,即一个存储体,不引发冲突
     int a = ldsMemory[index];
     /＊ 每个工作项访问8字节的LDS空间,即两个存储体,不引发冲突,
     且正好用满整个LDS的带宽 ＊/
     long l = ((long ＊)ldsMemory)[index];
     /＊ 每个工作项访问了16字节的LDS空间,即访问了4个存储体
     就拿前16个工作项而言,前8个工作项已经访问了8 ＊ 4 = 32个存储体
     而后8个工作项则访问了与前8个工作项同样的存储体。即,
     工作项0与工作项8访问了存储体0到存储体3,其中工作项0访问的是深度0,
     工作项8访问的是深度1
     工作项1与工作项9访问了存储体4到存储体7,其中工作项1访问的是深度0,
     工作项9访问的是深度1
     工作项7与工作项15访问了存储体28到存储体31,其中工作项7访问的是深度0,
     工作项15访问的是深度1
     因此,这引发了LDS存储体的两路冲突 ＊/
     float4 f4 = ((float4 ＊)ldsMemory)[index];
     /＊ 这就是最糟糕的LDS访问模式,工作项0到15都访问了LDS的存储体0。其中,工作项0访问了存储
体0的深度0;工作项1访问了存储体0的深度1,工作项14访问了存储体0的深度14;工作项15访问了存储体
0的深度15 ＊/
     int worst = ldsMemory[index ＊ 32];
     //这也不会引发存储体冲突,因为所有工作项都访问同一个LDS地址
     int stillOK = ldsMemory[0];