8.2 AMD RADEON R9 290X GPU

本节我们将讨论一下AMD GPU中的“波面阵”(在NVIDIA显卡中称为“线程块”)。这个概念是有助于我们区分OpenCL工作项和CUDA线程。不过,有时会无可避免的使用“线程”这个词来对GPU的工作方式进行描述。本书中我们将会使用“线程”来描述硬件线程。虽然,本节指定的设备为AMD Radeon R9 290X GPU,但是硬件线程与工作项映射的关系、线程调度和抢占的考量,以及内存系统的布局,不同厂商的同类产品其实很相似。

GPU上OpenCL的目标代码与CPU上的差别很大。读者应该还记得,图像处理器主要用来对3D图像进行渲染。这种设计让使得资源的优先级有所不同,与CPU架构也有很大的区别。目前GPU市场上,GPU的一些主要特性的差距越来越小,这些特性我们在第2章讨论过:

  1. 执行宽SIMD:大量执行单元对不同的数据,执行相同的指令。
  2. 大量多线程:GPU计算核支持并发大量线程。
  3. 硬件暂存式内存:物理内存可由编程者控制。

下面罗列出来的区别则更加微妙,不过其根据工作分配和交互对应用性能进行提高:

  • 支持硬件同步:并发线程间支持细粒度交互。
  • 硬件管理任务分配:工作队列管理与负载均和由硬件控制。

硬件同步能为计算单元中的多个波面阵减少同步开销,并且细粒度交互的开销非常低。

GPU为任务分发提供广阔的硬件空间,因为图像处理器支持三维图像。游戏任务在显卡的流水线上交错排布,让图形任务的管理变得更加复杂。图8.5中展示了OpenCL模型在AMD Radeon R9 290X GPU映射关系,其架构包含一个命令处理器和工作组生成器(生成工作组供硬件调度)。调度器会将计算任务分配到设备的44个核上。有关计算单元的内容将在8.2.2中详细讨论。

8.2 AMD RADEON R9 290X CPU - 图1

图8.5 OpenCL执行模型和内存模型在Radeon R9 290X上的映射

为了使用GPU获得高性能加速,调度需要十分高效。对于图像处理器,波面阵调度的开销必须很低,因为工作块可能很小(例如,一个像素只包含三个元素)。因此,当我们需要使用GPU获取较高的性能时,我们需要:

  • 提供大量的内核任务供硬件分发
  • 如果内核过小,考虑将内核进行合并

为每个内核提供足够工作,这样就能持续的占用工作组的流水线,所以调度器总能收到很多的任务,并且调度器通常会将给多的任务推到SIMD单元上。所以我们希望,设备每一次的吞吐可以创建大量的波面阵发送给GPU。

第二点是因为OpenCL设备支持队列模式。当OpenCL运行时选择执行相关设备上任务队列中的任务时,运行时会扫描队列中的任务,并选择合适大小的任务进行执行。这种任务由一个GPU工作命令缓存构成,其命令的执行放置在GPU的前端流水线上。其执行内容包括(1)构建一个队列,(2)将该队列放置在内存的某个位置上,(3)告诉设备这个队列在哪里,(4)要求设备执行这个队列。因为操作阶段这几个步骤是串行的关系,所以对于单个工作块来说,就会导致较高的延迟。另外,GPU在驱动层之下执行内核的过程中,为了能让GPU正常工作,需要在大量执行内核的上下文中进行切换。CPU端如果涉及到线程上下文切换,那应用的性能肯定不会太好。所以,如果任务量太小,就会出现多个上下文切换的情况,从而导致应用性能的下降。不过,这里有个开销常量,也就是对于单个工作队列的分配,以及之后对队列上任务执行的开销。通过传递大内核的方式(或较长的内核序列),可以避免这种开销的累积。这样做的目的,就是为了增加每个队列中所要处理的任务量。

8.2.1 线程和内存系统

CPU缓存层级结构的设计就是为了减少内存访问流的延迟,一旦内存访问流出现了重大的延迟,那么会对性能造成很大的影响。另外,因为GPU核使用线程和宽SIMD单元,所以其能在和标量相同的延迟开销下,将数据吞吐量最大化。图8.6展示了一个包含FX-8350 CPU和Radeon R9 290X GPU系统的内存层级结构。

这种内存系统中,GPU具有如下特征:

  • 大量的寄存器
  • 用于管理软件的便签式内存,在AMD硬件中称为共享本地数据(LDS, local data shares)
  • 高级别的片上多线程
  • 单个L2缓存
  • 高带宽内存

8.2 AMD RADEON R9 290X CPU - 图2

图8.6 分布式系统中的内存带宽

如之前所提到的,图像任务与计算任务有很大的区别,这就导致GPU需要有自己的执行和内存模型。实际上,GPU比CPU更加少的依赖数据重用,并且GPU的缓存大小要比CPU的小很多。即使在x86核和Radeon R9 290X计算单元上的L1缓存数据大小相同,GPU一次执行40个波面阵的话,每个波面阵所分配到的缓存更加的有限。GPU对于缓存的弱依赖是由一系列原因的,其中包括时域图像数据不可复用和由于数据集过大和大量的线程,无法保证为每个工作集提供足够的缓存空间。当数据复用发生时,多线程和高带宽内存将有助于克服缓存空间有限的情况。

GPU内存系统中,每个计算单元中有一个L2缓存。计算单元中的单L2缓存是为了保证GPU中L1缓存和外部缓存数据一致性而设计。如果向L2写入数据从而导致寄存器数据溢出,那么后续的每次访问则会有很高的的延迟,并且访存操作会阻塞在L2缓存处。

为了避免溢出,GPU提供了大量的寄存器。例如,x86架构下通用寄存器的数量是固定的(每个线程上下文中有16个通用寄存器),Radeon R9 290X显卡单个波面阵可用的寄存器数量为16384个!波面阵尝试只使用寄存器进行计算,并且也会使用到LDS,也会尽可能避免访存冲突。

LDS允许高带宽低延迟(编程可控)的读写访问。这种可编程模式的数据复用要比硬件控制更加高效。减少垃圾数据(加载到缓存而从不使用的数据)访问,可以认为LDS的容量要比同级的缓存小得多。另外,减少控制逻辑和结构体标记就可以节省出很多LDS空间。

GPU核中允许硬件控制多线程,可以很好的掩盖内存访问延迟。为了让性能和利用率达到最高,就要有足够多的波面阵执行起来。每个SIMD单元上能运行4个以上(包括4个)波面阵,或是很多应用需要每个计算单元上能运行16个波面阵。如果每个SIMD单元可以运行10个波面阵,那么就有显卡中就会同时有40个活动的波面阵。为了能更加快速的进行切换,波面阵的状态则保存在寄存器中,而非缓存中。虽然大量的波面阵可以掩盖内存访问延迟,但每个波面阵也都是需要资源支持的,这里必须要进行权衡。

为了减少每个波面阵产生的请求,系统中的缓存会通过一种过滤机制进行合并读取操作,将写入操作合并,尽可能一次性写入更多的内容——这种机制被称为合并访问(coalesing)。向量从规划好的内存(基于DRAM)中读取数据会更加的高效。

图8.6中CPU和GPU是通过PCIe总线连接在一起,两设备间的数据传输则由该总线完成。应为PCIe总线的带宽要比DRAM低很多,和片上内存的带宽没法进行比较。CPU和GPU交互比较多的应用中,带宽会成为应用的性能瓶颈。OpenCL应用中必须最小化与执行相关的数据拷贝。使用离散GPU时,如果应用总要将GPU执行的数据通过PCIe进行拷贝,那么应用就不可能达到最佳的性能。

8.2.2 指令集和执行单元

图8.7展示了基于AMD下一代核芯显示架构Radeon R9 290X中的计算单元。计算单元在该显卡中具有四个SIMD单元。当创建一个波面阵,其就会被指定到计算单元中的一个SIMD单元上(分配寄存器),并且在LDS上分配对应尺寸的内存。当波面阵有64个工作项,SIMD单元具有16个通道,那么SIMD单元上执行一个波面阵至少需要4个周期。这几个周期都需要一个SIMD单元上完成。这样的话,四个周期之后才能执行新的指令,确切的时间要从前一个指令完全提交到流水线上开始算起。

8.2 AMD RADEON R9 290X CPU - 图3

图8.7 Radeon R9 290X计算单元的微架构

回到计算单元,Radeon R9 290X除了SIMD单元,还有一定数量的执行单元。为了能够进行指令级别的并行,Radeon R9 290X计算单元不仅会将波面阵分配给SIMD单元,还会将对应的指令发送给其他硬件设备。调度器可能会将5个指令的每个周期放在不同单元上完成,这些单元包括SIMD单元、标量单元、内存单元、LDS或其他指定的功能硬件设备[2]。其他单元的调度与SIMD单元有些不同。比如,标量单元可以完成波面阵中的一条指令,其设计在每个周期可获得一个新指令。

之前提到的设备,比如第2章提到的Radeon HD 6970,其控制流就由分支单元全权管理。这种比较特殊的执行方式和市面上其他的向量结构有些不同。Radeon R9 290X的设计集成了标量和矢量,就像x86架构的CPU兼容SSE或AVX操作一样。

再回到有64个工作组的波面阵中,每个波面阵都具有一个指令流(64个工作项共用一个程序计数器),并且所有分支都以波面阵为粒度进行执行。为了在Radeon R9 290X上支持分支控制流,架构提供给了执行掩码,用以表示哪些部分写入或不写入结果。因此,波面阵中工作项间的任何分支,都需要使用架构指令集(instruction set architecture,ISA)对指令加以限制,将对应的操作的写入功能开启或关闭。这样执行单元必须指令块队列中的所有可能的分支覆盖到。这样的分支执行会让向量单元执行的效率大大下降。不过,这样的问题完全可以在编程阶段避免,编程者可以通过手动的方式向量化代码。相同的问题也会出现在其他具有计算能力的架构中,比如NVIDIA的GeForce GTX 780,其软件产品支持宽向量架构,同样也会遇到分支问题。应对分支问题的方式:编程者手动修改代码,编译器对代码进行调整,或是硬件向量化等。

下面的例子会运行在Radeon R9 290X的计算单元上(可见Radeon R9 290X海岛系列ISA标准[3])。下面运行一个简单的内核,在一维索引为0和不为0时所要执行的代码不同:

  1. kernel void foo(
  2. const global int *in,
  3. global int *out){
  4. if (get_global_id(0) == 0){
  5. out[get_global_id(0)] = in[get_global_id(0)];
  6. } else {
  7. out[get_global_id(0)] = 0;
  8. }
  9. }

这个简单的内核代码需要使用编译器进行编译,编译出的指令要与ISA进行映射。下面就让我们来看看,编译器是如何在Radeon R9 290X上翻译这段代码的。

  1. s_buffer_load_dword s0, s[4:7], 0x04
  2. s_buffer_load_dword s1, s[4:7], 0x18
  3. s_waitcnt lgkmcnt(0)
  4. s_min_u32 s0, s0, 0x0000ffff
  5. v_mov_b32 v1, s0
  6. v_mul_i32_i24 v1, s12, v1
  7. v_add_i32 v0, vcc, v0, v1
  8. v_add_i32 v0, vcc, s1, v0
  9. s_buffer_load_dword s0, s[8:11], 0x00
  10. s_buffer_load_dword s1, s[8:11], 0x04
  11. v_cmp_eq_i32 s[4:5], v0, 0
  12. s_and_saveexec_b64 s[4:5], s[4:5]
  13. v_lshlrev_b32 v1, 2, v0
  14. s_cbranch_execz label_0016
  15. s_waitcnt lgkmcnt(0)
  16. v_add_i32 v1, vcc, s0, v1
  17. s_load_dwordx4 s[8:11], s[2:3], 0x50
  18. s_waitcnt lgkmcnt(0)
  19. tbuffer_load_format_x v1, v1, s[8:11], 0 offen
  20. format:[BUF_DATA_FORMAT_32, BUF_NUM_FORMAT_FLOAT]
  21. label_0016:
  22. s_andn2_b64 exec, s[4:5], exec
  23. v_mov_b32 v1, 0
  24. s_mov_b64 exec, s[4:5]
  25. v_lshlrev_b32 v0, 2, v0
  26. s_waitcnt lgkmcnt(0)
  27. v_add_i32 v0, vcc, s1, v0
  28. s_load_dwordx4 s[0:3], s[2:3], 0x58
  29. s_waitcnt vmcnt(0) & lgkmcnt(0)
  30. tbuffer_store_format_x v1, v0, s[0:3], 0, offen
  31. format:[BUF_DATA_FORMAT_32, BUF_NUM_FORMAT_FLOAT]
  32. s_endpgm

清单8.1 Radeon R9 290X上分支操作所生成的ISA指令

这段代码表示在一个通道上所执行的所有操作(一个工作项)。不过,不同于高级语言,这里的标量操作(前缀为s_)会在图8.7所示的GPU标量单元上完成,矢量操作(前缀为v_)会在GPU矢量单元上完成。

清单中第11行(v_cmp_eq_i32 s[4:5], v0, 0),我们能看到这里执行的是矢量比较操作,将工作项ID(存与v0)与常量数0进行比较。比较的结果存于一个64位布尔掩码中,64位掩码由连续的两个标量寄存器组成(s[4:5])。得到的布尔掩码则用来决定OpenCL C内核所要执行的分支。

第12行(s_and_saveexec_b64 s[4:5], s[4:5]),这里隐式的将当前执行掩码与比较结果后的掩码进行与操作。例子中最后得到的执行掩码将用于判断执行if的哪个分支。目标标量寄存器中会将上次执行的掩码进行保存(本例中,s[4:5]将保存上次比较的结果掩码,目前存储的是上次的执行掩码)。当决定哪个工作项走else分支时,就需要使用上次得到的执行掩码。同样,在条件执行完成后需要上次的掩码用来重置执行掩码。这些操作保证对标量条件代码(SCC, scalar condition code)寄存器的设置,该寄存器用于触发条件分支。

设置SCC寄存器对程序来说是一种优化,因为其能保证对应工作项进入对应分支,不需要走的分支则不会去执行(本例中,当掩码为0,14行的s_cbranch_execz指令将保证if分支在执行过程中不会触发)。如果条件不满足,代码将进入else分支。本例中,工作项0将使if分支条件成立,其分支掩码为1,所以工作项0则会执行if分支中的语句——用一条SIMD通道完成(从第16行开始(v_add_i32 v1, vcc, s0, v1))。

当if分支执行完成,则会在第19行执行一个矢量加载(从tbuffer或纹理内存中加载数据,),将期望的数据存储到向量寄存器中,v1。第19行是if分支中最后一个操作。注意,在OpenCL C代码中在加载操作完成后,还有一个存储操作。因为,不同分支所要写入的地址是相同的,所以这里编译器将这个存储操作放在程序的最后进行(tbuffer_store_format_x v1, v0, s[0:3], 0, offen)。

第22行(s_andn2_b64 exec, s[4:5], exec)会将当前执行掩码与原始执行掩码进行与操作。新的执行掩码则表示当前工作项执行的是else分支的语句。else分支不需要从内存中加载数据,工作项只需要将常量0存储v1即可(第23行(v_mov_b32 v1, 0))。注意不同的分支都将结果保存在v1寄存器中,这样就可以将存储操作放在之后完成。

细心的读者可能已经发现了,指令中这里并没有对else分支进行忽略。本例中,编译器认为与其忽略else分支,不如不对其进行加载,这样就节省了执行预测的性能开销。而且,这样操作就不会对v1寄存器进行更新,和忽略else分支达到了同样的效果。

这是一个简单的例子。在当有嵌套或深层条件时,就会有更加复杂和更加长的序列用来存储条件掩码,并且与执行掩码进行与操作,逐级缩小执行管道最后所要执行的变量分支。逐级缩小的过程将会对造成执行性能的下降。所以,良好的代码设计将会在执行相同指令的前提下,获得更佳的性能。程序的掩码管理十分复杂,并且GPU向量操作与CPU向量操作(SSE)也区别。其在多核系统下,向量操作就并非是一种抽象的概念了。

最后,第24行(s_mov_b64 exec, s[4:5]),将执行掩码进行重置,并且将数据存储在v1寄存器中,在第30行(tbuffer_store_format_x v1, v0, s[0:3], 0, offen)使用tbuffer将数据存储到对应位置上。

8.2.3 资源分配

GPU中的每个SIMD单元都有固定数量的寄存器和LDS存储空间。每个计算单元中有256KB的空间供寄存器使用。这些寄存器分为四组,每个SIMD单元有256个寄存器,每个寄存器,每个单元有64个通道,每个通道具有32位。计算单元中,寄存器是执行波面阵的最小单位。每个计算单元有64KB的LDS,就像对32行的SRAM一样可以进行随机访问。LDS会在计算单元中分成多份,划分给对应的工作组,用于开辟对应的局部内存(通过OpenCL运行时传入的参数进行分配)。

计算单元执行内核时,我们可能看到如图8.8所示的资源瓶颈。图中两个波面阵,每个波面阵中有两个工作组,每个工作项需要42个向量寄存器(按照波面阵的尺寸等比例放大),并且共享50个标量寄存器,并且每个工作组需要24KB LDS。每个计算单元能够同时执行4个这样的波面阵,这样可将每个波面阵分配给一个SIMD单元,同时运行4个会让设备保持忙碌状态。不过,当执行标量代码或内存操作时,就没有空间进行切换操作。下面给出在Radeon R9 290X上工作组对一个计算单元的占用比(等式8.1):

8.2 AMD RADEON R9 290X CPU - 图4 (8.1)

8.2 AMD RADEON R9 290X CPU - 图5

图8.8 R9 290X GPU物理内存与OpenCL内存间的映射关系

等式8.1中,VGRP表示矢量通用寄存器,SGPR表示标量通用寄存器,WI表示工作项,WG表示工作组,WF表示波面阵。通过OpenCL内核参数传入的部分使用粗体表示。

图8.8的例子中,如果增加SIMD单元中波面阵运行的数量(4个或更多),在控制流和内存延迟方面,让标量和矢量单元忙碌起来,有助于对延迟的隐藏。本例中LDS有限,对于算法来说,将波面阵数量(只有一个波面阵)增加到3个是个不错的选择。另外,减少LDS内存分配,则可让每个工作单元多执行三分之一个工作组。这个方式在波面阵等待栅栏或内存访问时非常有用(SIMD单元闲置状态)。

每个波面阵运行在一个SIMD单元上,直到执行完成。任意组具有相同工作组的波面阵都会在同一个计算单元上完成。由于每个工作组的状态存储量的原因,才将这些工作放在同一个计算单元完成。本例中,我们看到每个工作组需要24KB LDS,且寄存器可操作的空间超过21KB。在刷新内存和搬移到其他核上时,这个数据量还是相当大的。当内存控制器执行高延迟的读或写操作时,如果没有其他波面阵在计算逻辑单元(ALU)中时,那就没有任务分配到SIMD单元上,计算硬件这时处于闲置状态。


[2] Advanced Micro Devices, GCN Whitepaper. AMD Graphics Core Next(GCN) Architecture, 2013. http://www.amd.com/Documents/GCN_Architecture_whitepaper.pdf

[3] Advanced Micro Devices, ADM Sea Islands Series Instruction Set Architecture, 2013. http://developer.amd.com/wordpress/media/2013/07/AMD_Sea_Islands_Instruction_Set_Architecture.pdf.