nVidia GPGPU vs AMD Radeon HD Graphics执行模式对比

大家做高性能计算的朋友，想必对CPU的执行模式已经非常熟悉了吧。当代高级些的CPU一般采用超标量流水线，使得毗邻几条相互独立的指令能够并行执行——这称为指令集并行（ILP，Instruction-Level Parallelism）；而像x86引入的SSE（Streaming SIMD Extension）、AVX（Advanced Vector Extension），以及ARM的NEON技术都属于数据级并行（Data-Level Parallelism）。而GPGPU的执行与CPU比起来还是有不少差异的。这里，为了能够让大家更好地理解、并使用OpenCL，想谈谈当前主流用于超算的GPGPU的执行模式。

下面主要针对nVidia的Fermi和Kepler架构以及AMD的TeraScale3（Radeon HD 6900系列）和GCN架构进行分析。

我们先来简单介绍一下OpenCL中的一些术语对应到nVidia以及AMD的GPGPU硬件中的称谓。

在物理上，一个GPGPU作为一个计算设备，在OpenCL中就称为Device。在这么一个计算设备中由若干大的计算核心构成，这个大的计算核心在OpenCL中称为CU（Compute Unit），在nVidia中称为SM（Streaming Multiprocessor），在AMD中则称为SIMD（Single-Instruction-Multiple-Data）。一个大的计算核心中又由许多小的计算核心构成，这种小的计算核心在OpenCL中称为PE（Processing Element），在nVidia和AMD中均称为SP（Stream Processor）。此外，nVidia也好、AMD也罢，甚至还有一些如Intel HD Graphics这样的GPGPU，有一个OpenCL中没有对应到的术语，它其实属于一种GPGPU的执行模式，我这里暂且称为“线程流”。以应用开发者的角度来看，它是GPGPU中线程执行的最小并行粒度，稍后会详细讲解。这个概念在nVidia中称为Warp，在AMD中称为Wavefront。

我们下面先来看看nVidia的Fermi架构是如何在OpenCL中执行的。在Fermi架构中，一个SM一共有32个SP，16个存储器读写单元，四个特殊功能计算单元SFU（用于计算超越函数等复杂操作），64KB的共享存储器（Local Memory），32768个32位寄存器，两个Warp调度器与两个指令分派单元。

在此架构中，SM调度线程时会将32个线程（OpenCL中称为work-item）组成一组，然后并发执行。这个32线程组就是一个warp。由于每个SM含有两个warp调度器与两个指令分派单元，因此这就能够将两个warp同时进行发射和执行。Fermi的双warp调度器先选择两个warp，然后从每个warp发射一条指令到一个十六核心的组，或是十六个读写单元，或是四个SFU。正由于warp执行是独立的，因此Fermi的调度器无需检查指令流内的依赖性。

那么这个执行模式如何映射到一个OpenCL的kernel程序里呢？我们现在假设给kernel分配了512个work-item（Fermi架构的Max Work-group Size为1024），work-group size也是512，然后执行以下kernel代码：

 kernel ocl_test(__global int *p)
 2 {
   
   
 3     int index = get_global_id(0);
 4 
 5     int x