OpenCL AMD GCN架构的GPU

在实际的应用中,编译器很难抽取到足够的VLIW并行度来满足Cayman的需求,而很多程序员并不愿意手动地在OpenCL代码中引入SIMD操作,再加上在一些代码中引入SIMD操作需要处理许多烦心的边界问题,因此在实际的很多代码中,Cayman架构的效率偏低。为了解决这一缺陷,为程序员提供更好的编程平台,AMD推出了GCN架构。

GCN整体架构与执行模型

由于VLIW执行引擎架构后面应该不会再有了。现在主流GPGPU的设计都是采用标量处理器元素(PE)的执行模式。也就是说,在OpenCL设备查询中我们看到数据操作类型偏好为int、f loat等标量类型,而不是像int4、f loat4等向量类型。然而由于GCN架构设计的灵活性,无论我们采用向量数据类型还是标量数据类型都能达到性能峰值。也就是说当我们用OpenCL编程时,即便基于以前老的VLIW4的模型,在GCN(Graphics Core Next)架构的GPGPU上运行,性能也会表现出色。
另外,在AMD的官方命名上,GCN架构的硬件单元术语与OpenCL模型中的术语用法完全一致,如计算单元(Compute Unit,CU)、处理元素(Processing Element,PE)等。下面会做详细介绍。
下面我们看一款GCN架构处理器的简单结构体,下图是AMD Radeon HD 79XX系的GPU。

GCN架构图

从上图中我们可以看到,AMD Radeon HD 79XX系的GPU含有两大计算阵列,每个计算阵列含有一个计算引擎/命令处理器,16个CU。两个阵列共享一个L2 Cache。而每个CU含有一个标量单元,4个向量单元,一个L1数据Cache和一块LDS。下面我们把CU放大看看其内部结构。
下图是AMD GCN架构CU的内部结构示意图,下面详细解说。



GCN CU内部示意图

从上图我们可以看到,在GCN架构中,每个CU含有一个标量单元用于做分支、同步等操作;4个独立的向量单元(SIMD0、SIMD1、SIMD2、SIMD3)用于做主要的向量算术逻辑计算处理。这4个SIMD单元的每一个可以同时对16个工作项执行一单次操作。而每个SIMD单元每一次只能在它自己的wavefront上执行。在GCN架构中,其wavefront与Cayman架构类似,对应64个工作项。执行完整条wavefront需要花费4个周期,每个周期执行其中连续的16个工作项。
我们看看一个CU的执行前端。在GCN架构中,每个SIMD具有其自己独立的40位程序计数器,10条wavefront以及用于执行这些wavefront的指令缓存。因此,整个CU就含有40个wavefront。每条wavefront可以在不同的工作组或是不同的kernel中执行。对于刚才的Radeon HD Graphics 79XX系的GPU而言,它一共含有32个CU,而每个CU包含40条wavefront,每个wavefront正好对应了64个工作项,因此我们可以在某一时刻看到它对32×40×64=81920个工作项进行操作。
在GCN中,4个CU可以合并成一个簇(cluster)共享一个32KB的四路组相联的L1指令Cache。Cache行长度为64字节,一般可以保留8条指令(每条指令8个字节)。当Cache满的时候有一个新的请求到来时,Cache控制器将会使用最近最少使用策略将旧的指令逐出Cache,为新的指令腾出空间。被4个CU所共享的L1指令Cache具有4个存储体(其实每个存储体就是四路组相联的一路),可以在每个周期对所有这4个CU持续地取32字节的指令(如果像上述所提到的每条指令为8个字节,那么一共就正好取了4条指令)。而在一个CU内,取指令会在4个SIMD单元之间基于执行时长、调度优先级以及wavefront指令缓存的利用率等因素进行调度。
一旦其中一个SIMD单元将指令取到其wavefront缓存中之后,下一步就是译码并发射指令。这里再提一下,CU在每个周期会使用时间片轮询方式(并且基于一些调度策略)选择其中一个SIMD单元进行译码并发射一条指令。被选中的SIMD然后就可以从自己10条wavefront的指令缓存译码并发射5条指令到执行单元。每个CU具有16个缓存用于跟踪栅栏指令,栅栏指令用于强制将一条wavefront做全局同步(在当前的工作组内)。
上面我们看到SIMD单元在一个周期可以并行发射5条指令,但这5条指令若要并行发射是有条件的。
第一个约束是,这5条指令必须是不同指令类型的。CU前端可以译码并发射的指令类型有7种,分别是:分支指令、标量ALU或标量访存、向量ALU、向量访存、LDS访问操作、全局数据存储器的访问、其他特殊指令(诸如栅栏等同步指令)。对于每个SIMD单元,每个周期每种类型只能有一条指令被发射,这是为了避免执行流水线被过度消耗。
第二个约束是,这5条指令中,每一条都必须来自不同的wavefront(之前已经提到一个SIMD单元含有10条wavefront),这是为了维持指令执行次序。除了这两个约束外,没有其他限制。因此,这其实给了编译器以及开发者很大的自由度来安排指令执行。同时,我们也可以通过并发执行多个内核程序来充分利用GCN架构GPU的指令并行执行功能。
一个CU中的每个SIMD向量单元含有512个标量寄存器(SGRP)用于处理分支、常量以及其他数据常量,贯穿整条wavefront。并且还含有256个向量寄存器(VGPR)用于算术逻辑计算。向量寄存器实际上也是标量的,不过它们会贯穿整个wavefront被复制给该wavefront中所有工作项使用。
以上我们已经详细地介绍了GCN架构GPU的整体结构以及执行特性。下面,我们将给出一个图来体现一个CU中的4个SIMD单元是如何并行执行一条指令的(这里不考虑指令并行执行的特性)。由于每个SIMD单元其实只含有一个标量的ALU(算术逻辑单元),而可以针对16个工作项进行操作。并且每个周期,CU会挑选一个SIMD单元把指令送到其wavefront指令缓存中。所以,一个CU中的4个SIMD单元其实是以一种流水的方式并行执行的,如下图所示。

SIMD单元执行wavefront示意图

对于上图的执行过程,我们假设当前就只有向量算术逻辑指令在执行,因此每个SIMD单元就执行一条指令。一开始,假定当前CU的4个SIMD单元的wavefront都处于空闲状态,在第0个周期,该CU的SIMD0先被选中,然后执行该指令操作;在第1个周期,由于SIMD0此时的wavefront已经处于活跃状态,所以调度器选中了SIMD1开始译码并发射指令;在第2个周期,SIMD0与SIMD1都处于活跃状态,SIMD2被选中,然后开始译码并发射指令进行执行;在第3个周期,SIMD0、SIMD1和SIMD2都处于活跃状态,因此选中SIMD3译码并发射指令开始执行。而到了第4个周期,由于SIMD0已经完成了整条wavefront,因此如果存在更多的工作项(比如给当前内核分配了512个工作项),那么该CU再次调度到SIMD0开始执行其第二条wavefront。为了避免歧义,这里补充说明一下,所谓的第二条wavefront也有可能是之前执行的那条,根据当前该SIMD的10条wavefront的活动状态进行选择。图8-8所示的执行流程是对一个256个工作项的工作组进行执行的过程。这里再强调一下,任一wavefront都会执行连续的64个工作项,因此不会存在一条wavefront在第1个周期执行了工作项0到15,而第2个周期又去执行工作项64到79的情况。

GCN架构GPU的LDS结构与优化

GCN架构GPU的LDS结构与Cayman架构的类似,不过在某些细节上有些不同。在GCN架构中,LDS比L1 Cache访问速度都快很多,它具有L1 Cache双倍的峰值带宽,更低的延迟,以及对原子操作的高性能支持。因此,对于产生中间结果且需要被多次重用的数据,将它们合理地安排在LDS中并合理地进行访问能给内核程序的执行带来很好的性能收益。GCN架构的LDS另一个特点是,它不需要合并(coalescing)访问。所谓合并访问,简单来说就是对于1/4wavefront的工作项,每个工作项索引ID访问其对应的存储地址。比如:int index=get_global_id(0);int a=mem_buffer[index];这种对连续存储地址的访问一般对于全局存储器的访问来说能带来非常高的性能。而对于LDS访问来说,即便存在交叉访问的情况(只要不产生存储体冲突),仍然能达到带宽峰值。比如:int index = get_global_id(0);int a=mem_buffer[index^1];。
这里想提醒各位读者的是,所有AMD Southern Islands、Sea Islands以及Volcanic Islands的GPU(这三种型号都基于GCN架构)每个CU含有64KB的LDS,但是只有其中的32KB分配给每个工作组。也就是说,应用开发人员只能在内核程序中为每个工作组分配最多32KB的局部存储器空间。这意味着软件开发人员需要为每个CU准备至少两个可同时调度的工作组才有可能完全利用LDS。由于一个内核程序对寄存器的分配完全根据编译器在编译时候的计算得出的,而不是通过运行时来控制。编译器能够得知当前的GPU架构,但无法知道主机端给每个工作组分派了多少个工作项,一共有多少个工作组。然而,利用已知架构来分派额外的寄存器使用是可能的。另外,如前所述,一个CU能调度多个工作组,具体能调度多少个工作组需要看当前执行资源,其中LDS的使用也是其中一个指标。GCN的LDS在结构上与Cayman的差不多,一共含有32个存储体,每个存储体是4字节宽,256字节的深度。GCN架构的LDS会以半wavefront(即连续的32个工作项,即一个工作组中的工作项0到工作项31;工作项32到工作项63)为单位做所访问地址的存储体冲突检查,而不是每个周期检查一次。
像AMD Radeon HD 7970 GPU,其局部存储器可以在一个周期内为每个存储体服务一个请求(即每个周期可以做高达32个存储体的访问)。其全局存储器的带宽在100GB/s左右,而局部存储器带宽可高达3.5TB/s,是全局存储器带宽的几十倍之多。然而,如果一个工作组内的所有工作项都访问同一个存储体,那么这些访问将会以连续的周期做串行处理。LDS操作本身不延迟;然而,编译器会在发射依赖于之前计算结果的操作之前插入一个等待操作。如果我们对一个局部存储器读数据并且产生了存储体冲突,但后续不急着使用,而是做其他计算,等这些计算做好之后再用栅栏操作做一次同步的话,那么局部存储器的存储体冲突在内核程序执行过程中不会被体现出来,它们会被计算执行给掩盖掉。如果内核程序的计算比较简单,同时又对于局部存储器的访问模式会产生大量存储体冲突,此时使用常量存储器或图像存储器可能会比局部存储器要高效些。
因此,对局部存储器LDS高效使用的关键在于对其访问模式的控制。在同一周期对LDS不同存储体的访问,尤其是由1/4的wavefront所发出的LDS访存请求要避免存储体冲突。与Cayman架构一样,如果1/4的wavefront中每个工作项访问LDS对应的相邻两个存储体(即每个工作组对LDS相应位置读写8字节数据)会产生对LDS访问的峰值,而不会造成存储体冲突。而如果每个工作项一次读写16字节数据(比如使用f loat4进行读写),那么会至少产生两路存储体冲突。

GCN架构GPU的全局存储器与优化

GCN架构的GPU,每个CU含有16KB大小的L1 Cache。L1 Cache的带宽峰值为:CU个数×(4线程/周期)×(16字节/线程)×引擎时钟周期。
如果两个存储器访问被导向了同一个控制器,那么硬件会对这两个访问做串行处理。这被称为一次通道冲突(channel conf lict)。类似地,如果两个对存储器访问请求走向了同一个存储体,那么硬件也会对这两个访问做串行处理。这被称为一个存储体冲突(bank conf lict)。就应用开发者角度而言,通道冲突与存储体冲突没有什么太大不同。一个巨大的2的幂次跨度访问往往会导致一次通道冲突。导致通道冲突的2的幂次跨度的具体大小取决于芯片上冲突依赖的类型。对于具有8个通道所导致的一次通道冲突的跨度可能会导致具有4个通道机器上的一次存储体冲突。
最重要的概念就是存储器跨度:在存储器地址中的递增是以元素为单位进行测量的,也就是在一个内核中由两个相邻工作项所读或写的相继两个元素之间的跨度。许多重要的内核并不专有地使用单一种跨度访问模式;而是它们具有许多非均匀的跨度访问模式。比如,许多代码在二维或三维数组的每一维度执行类似的操作。在低维度上执行计算经常用单位跨度(即相邻两个工作项之间访问存储器位置的跨度)来完成,而在其他维度上的计算跨度一般都会比较大。这时,如果代码毫无改动就搬到GPU上执行时就会造成严重的性能下降。而具有Cache的CPU也会存在相同问题,对存储器访问采用巨大的2的幂次跨度来访问数据只能将一小部分数据放入Cache行中。
一种解决方法是重写代码,在执行两个内核命令之间对数组做个变换。这允许所有计算都以单位跨度完成。当然,这里必须确保数组变换所花费的时间与执行内核计算比较起来足够小。
对于含有多个内核的情况下,性能下降会十分厉害,因此尝试理解并解决这个问题是很有价值的。
在GPU编程中,最好的模式是让相邻的两个工作项读写相邻的存储器地址。这是一种有效避免通道冲突的方式。
当应用程序完成对访问模式与地址生成的控制时,开发者必须将数据结构小心安排好,以最小化存储体冲突。每个工作项对不同的存储体进行访问是可并行操作的。如果含有存储体冲突,那么这些造成冲突的访问只能被串行执行以解决冲突。
看下面这个比较极端的例子:

for(char *ptr = base; ptr 〈 max; ptr += 16 * 1024)
r0 = *ptr;

这里,循环中每次迭代对ptr地址空间的访问都是跨一个L1 Cache尺寸的跨度,对存储器所请求的访问都是同一个通道上的同一个存储体,所以所有请求都会被串行处理。
这是需要被避免的低性能模式。当跨度为2的幂次(比通道间隔的跨度大)时,上述的循环只能访问存储器的一个通道。
在所有AMD Radeon HD 79XX系列GPU上,一共有12个通道。一条crossbar将加载分发到合适的存储器通道。每条存储器通道具有一个读/写全局L2 Cache,分给每条通道64KB,Cache行为64字节。
在AMD Radeon HD 78XX GPU上,一个以每256字节来切换通道。由于wavefront为64,所以如果一个wavefront中的每个工作项从一个64字(每个字长度为4字节)区域读不同的地址,那么就能避免通道冲突。所有AMD Radeon HD 7XXX系列GPU都具有相同的通道布局。
在AMD Radeon HD 7XXX系列GPU上,在每个CU上能执行一个或多个工作组。工作组以一个线性次序进行分发,当然分发速度很快。对于一维,就好比:分发次序= get_group_id(0);对于二维,就好比:分发次序 = get_group_id(0)+ get_group_id(1)* get_num_groups(0)。这是索引空间中的以行为主的块。一旦所有CU在使用中时,那么其他尚未被CU所使用的工作组会根据需要(即根据当前CU可用资源以及当前活跃的工作组执行被隐退)分配给CU。工作组按次序隐退,因此活跃的工作组是连续的。在6.7节中也提到了一般GPU调度工作组的特性。如果当前所有活跃的工作组的执行依赖于非活跃的工作组,那么会造成死等而使得执行超时。
在任一时刻,每个CU从一条wavefront执行一条指令。在访存密集型内核程序中,就好似该指令是一个访存操作。由于在79XX GPU上一共有12条通道,所以在一个周期内最多只能有12个CU发射一条存储器访问操作。如果12条wavefront的访存都走向不同的通道,那么此时是最最高效的。要达成这个目的的一种方式是为每个wavefront访问256字节(64个工作项,每个工作项访问4字节)的连续组。
对于Southern Island设备尽管不支持合并写,不过如果在一个工作组内做连续的地址访问仍然能提供最大性能。每个CU以1/4的wavefront为单元访问存储器系统。CU为每个工作项传输一个32位地址以及一个元素大小的数据。这产生了每个1/4的wavefront一共16个元素+16个地址。在GCN架构的设备上,在数据被传送到存储器控制器之前需要两个周期来处理1/4的wavefront。

赞(2)
未经允许不得转载:极客笔记 » OpenCL AMD GCN架构的GPU
分享到: 更多 (0)

评论 抢沙发

  • 昵称 (必填)
  • 邮箱 (必填)
  • 网址