NVIDIA's GT200--- Inside a Parallel Processor

本文是Davick kanter博士的《NVIDIA's GT200--- Inside a Parallel Processor》的部分中文译文


简介

过去十年中，计算领域出现了一个新的发展趋势。

由Intel,IBM,SUN,AMD和富士通生产的通用CPU虽然有了很大发展，但性能提高速度却已经不能与与上世纪八十年代末九十年代初相比。单线程处理性能在很大程度上受到了限制。这些限制一方面来自于通用计算程序中过低的指令级并行;另一方面来自于“功率墙(Power Wall)”--集成电路的功率消耗的物理限制。摩尔定律能够为处理器提供数以亿计的晶体管，但在为了运行单线程程序而设计的处理器中，这些晶体管绝大多数都被用于制造高速缓存(Cache)。这样做虽然能把处理器功耗控制在合理的范围内，却阻碍了性能的进一步提高。

与此同时，GPU（图形处理单元）却有效的利用了数量巨大的晶体管资源。由于图形渲染过程的高度并行性，GPU的性能按时间呈几何级数增长。同时，随着GPU计算能力的不断增长，一场GPU革命的时机也成熟了。GPU已经从由若干专用的固定功能单元(Fixed Function Unit)组成的专用并行处理器，进化为了以通用计算资源为主，固定功能单元为辅的架构。

这一进化过程的最初征兆出现在以ATI R300和Nvidia NV30位代表的DirectX 9时代的GPU上。与前代的DirectX 8时代GPU相比，这一代的显卡拥有可编程的像素和顶点渲染器，具有了一定的浮点计算能力。AMD收购ATI也是一个微妙的转折点—大多数人当时还在关注AMD与Intel在市场上的竞争，却没有意识到AMD的收购行为有着更加复杂的动机。

支持DirectX 10的显卡的首次出现是一个分水岭：过去只能处理像素或者只能处理顶点的专门功能处理单元被通用的统一着色器架构(Unified Shader Architecture)取代了。以ATI R600和Nvidia G80为代表的DirectX10时代GPU能够提供了超越以往任何GPU的能力：它们拥有数百个功能单元，能够处理一部份过去只能在CPU上运行的数据并行问题。需要强调的是，这些第一代的DirectX 10 GPU只能处理“一部分”数据并行问题；它们只适合处理使用类似阵列的数据结构，并且高度并行的问题。这些处理器的双精度浮点的计算能力十分有限；它们虽然基本符合IEEE规范对32位单精度浮点的大多数规定，但同时也缺少对异常的处理，并省略了一些舍入模式。

这场革命的结果使计算领域的局势突然变得复杂了许多。现在不仅有各种各样的CPU，也有了能够处理并行计算任务的GPU。这些产品的计算能力、可编程性和适合进行的计算任务各有不同。GPU、Cell和Niagara这一类的产品优缺点都十分突出：在处理需要复杂分支的单线程任务时，它们的表现简直无药可救；而在处理充分并行的任务时，它们的性能比CPU能够强上数十甚至数百倍。在可编程性上，Niagara和通用CPU的编程较为灵活，GPU很难处理复杂的数据结构，而Cell简直和程序员有仇。

讽刺的是，虽然ATI与CPU厂商AMD合并了，但另一家GPU厂商Nvidia却在通用计算方面有更加全面而坚实的基础。本文主要介绍Nvidia GPU在通用计算领域方面的应用，尤其是CUDA和最新一代GT200 GPU。GT200架构被使用在了面向娱乐市场的GeForce，用于高性能计算的Tesla和用于专业渲染领域的Quadro产品线上。本文并不深究DirectX 10和OpenGL2.1中错综复杂的现代3D流水线，只有在与CUDA的编程模型进行类比时才会提到。



CUDA 执行模型

Nvidia的并行编程模型被命名为CUDA（Computing Unified Device Architecture，统一计算架构模型）。CUDA的基本思想是尽量得开发线程级并行(Thread Level Parallel)，这些线程能够在硬件中被动态的调度和执行。CUDA编程模型的重点是将CPU做为终端(Host)，而GPU做为服务器(Server)或协处理器(Coprocessor)，或者设备（Device），从而让GPU来运行一些能够被高度线程化的程序。所以，GPU只有在计算高度数据并行任务时才能发挥作用。在这类任务中，需要处理大量的数据，数据的储存形式类似于规则的网格，而对这写数据的进行的处理则基本相同。这类数据并行问题的经典例子有：图像处理，物理模型模拟（如计算流体力学），工程和金融模拟与分析，搜索，排序。而需要复杂数据结构的计算如树，相关矩阵，链表，空间细分结构等，则不适用于使用GPU进行计算。找到程序中的计算并行度后，就能将一部分程序移植到GPU上。运行在GPU上的程序被称为Kernel(核)。核并不是完整的程序，而是整个程序中的若干基本的关键数据并行计算步骤。

图1 CPU与GPU 串行程序，核，网格和线程块

除了数据并行的核以外，程序中也有标准的串行程序。如表一所示，在两个核之间运行的就是串行代码（如果两个核函数之间没有串行代码，它们就可以被合并为一个）。理论上，串行代码的作用只是清理上个核，启动下一个核。但由于目前的GPU的功能仍然十分有限，串行部分的工作仍然十分可观。

核以网格(Grid)的形式执行，每个网格由若干个线程块（block）组成，每一个线程块又由最多512个线程(thread)组成。属于同一线程块的线程拥有相同的指令地址，能够并行执行，并且能够通过共享存储器(Shared memory)和同步栅（barrier）进行线程块内通信。同一线block中的thread开始于相同的指令地址，理论上能够以不同的分支执行。但实际上，在block内的分支因为性能方面的原因被大大限制了。

核实质上是以block的形式执行的。CUDA引入了grid这个概念来表示一系列能够（并不一定）完全并行执行的线程块的集合。线程块的执行没有顺序，完全并行。这是CUDA的一个很杰出的特性：无论是在一次只能处理一个线程块的GPU上,还是在一次能处理数十乃至上百个线程块的GPU上，这一模型都能很好的适用。

目前，一个核函数只对应一个网格，但是未来这一限制极有可能会被解除。实际使用中在一个网格访问数据的同时，如果能够进行另一个网格的计算，则可以有效的提高设备的利用率。但现在网格以串行代码划分了边界，不能很好的隐藏访存延迟(Memory Access Latency)。

因为同一线程块中的线程需要共享数据，因此它们必须在同一个处理器（Nvidia称之为Streaming Mulitporcessor，流多处理器，缩写SM）中发射。线程块中的每一个线程被发射到一个执行单元（Nvidia称之为Streaming Processor，流处理器，缩写SP）。这里涉及到了Nvidia GPU的内部架构，我们将在下一节进行详细说明。

CUDA存储器模型除了执行模型以外，CUDA也规定了存储器模型（如图2所示）和一系列用于主控CPU与GPU间通信的不同地址空间。图中红色的区域表示GPU片内的高速存储器，橙色区域表示DRAM中的的地址空间。
首先，是最底层的寄存器(register,REG)。对每个线程来说，寄存器都是私有的--这与CPU中一样。如果寄存器被消耗完，数据将被存储在本地存储器(local memory)。本地存储器对每个线程也是私有的，但是数据时被保存在帧缓冲区DRAM中，而不是片内的寄存器或者缓存中。线程的输入和中间输出变量将被保存在寄存器或者本地存储器中。
然后是用于线程间通信的共享存储器。共享存储器是一块可以被同一block中的所有thread（上节提到过，一个block最多可以有512个thread）访问的可读写存储器。访问共享存储器几乎和访问寄存器一样快，是实现线程间通信的延迟最小的方法。共享存储器可以实现许多不同的功能，如用于保存共用的计数器（例如计算循环迭代次数）或者block的公用结果（例如计算512个数的平均值，并用于以后的计算）。
除此以外，还有两种只读的地址空间—常数存储器和纹理存储器（constant memory and texture memory），它们是利用GPU用于图形计算的专用单元发展而来的。常数存储器空间较小(只有64k)，支持随机访问。纹理存储器尺寸则大得多，并且支持二维寻址（传统的缓存只支持一维寻址）。
这两种存储器实际存在于帧缓冲区DRAM中，但由于它们的只读性质，在GPU片内可以进行缓存，从而加快访问速度。这两种存储器并不要求缓存一致性—它们是只读的。但这也意味着如果CPU或者GPU要更改常数存储器或者纹理存储器的值，缓存中的值在更新完成之前也无法使用。CUDA程序中，常数存储器用于存储需要经常访问的只读参数，而是用插值或者滤波的纹理存储器访问对大尺寸的二维或者三维图象或者采样序列进行高带宽的流式访问。
最后是全局存储器（global memory），使用的是普通的显存。整个网格中的任意线程都能读写全局存储器的任意位置，并且既可以从CPU访问，也可以从CPU访问。由于全局存储器是可写的，GPU片内没有对其进行缓存。

CUDA 编程接口（API）

CUDA的API由一系列软件组成。处于其核心地位的是CUDA C语言，可以由Nvidia的CUDA编译器（nvcc）编译。值得指出的是，CUDA C不是C语言，而是对C语言进行扩展的变种。CUDA对C的扩展主要包括以下四个方面：

对函数进行标示，指明函数是在CPU上运行，还是在GPU上运行
对变量进行标示，指明变量存储在GPU的哪一个地址空间上
通过设置网格和线程块，控制核函数的并行执行
为了区分不同的线程块和线程，加入了内建变量
CUDA API需要CUDA驱动才能运行，新版本的Nvidia驱动已经包含了CUDA驱动。CUDA运行时是CUDA API的可选组件，它是一种动态编译器（JIT），能够直接访问实际中的底层硬件架构。最后，API还包括数学函数库（math library），cuFFT，cuBLAS和cuDPP。用户可以调用这些库函数，快速的开发显卡程序。

 

Nvcc编译器可以生成三种不同的输出：PTX，CUDA二进制序列和标准C。PTX(Parallel Thread eXecution)作用类似于汇编，是为动态编译器（包含在标准的Nvidia 驱动中）设计的输入指令序列。这样，不同的显卡使用不同的机器语言，而动态编译器却可以运行相同的PTX。这样做使PTX成为了一个稳定的接口，带来了很多好处：后向兼容性，更长的寿命，更好的可扩展性和更高的性能，但在一定程度上也限制了工程上的自由发挥。这种技术保证了兼容型，但也使新一代的产品必须拥有上代产品的所有能力，这样才能让今天的PTX在未来的系统上仍然可以运行。



虽然PTX和JIT编译器提供了很高的性能，但也不是在所有的场合都适用。某些独立软件开发商倾向于牺牲性能，以获得更好的可确定性和可验证性。JIT编译器的输出随着目标硬件和一些其他因素会发生变化。对于需要能够确定的代码的独立软件开发商（比如很多财经软件开发商），它们可以将代码直接编译成CUDA二进制代码，这样就能避免JIT过程的不确定性。直接编译得到的CUDA二进制代码是与特定的硬件和驱动相关的。



Nvcc的输出还包括标准C。由nvcc生成的C代码将被重定向到其他编译器进行编译，比如ICC，GCC或者其他合适的高性能编译器。CUDA中明确的表示了程序中的并行度没不仅在用于编写运行在Nvidia GPU上的代码时非常有效，而且为多核CPU生成高性能代码。在某些应用中，CUDA生成的代码比标准的x86编译器生成的代码的性能提高了4倍。



Evolution of the ISA and Compute Capabilities



CUDA的版本随着底层硬件的进步也在演化。CUDA计算设备（Nvidia GPU）的计算能力以版本号表示：版本号的第一位用于表示核心架构，第二位（小数点后）则表示了更加细微的进步。



虽然CUDA硬件只诞生了短短一年多的时间，但已经进行了三次重大版本更新，在通用计算能力上有了很大的提高。



版本1.1增加了32bit字在全局存储器中的原子操作(Atomic Operation)功能；



版本1.2增加了32bit字在共享存储器中的原子操作功能，以及64bit字在全局存储器中的原子操作功能，增加了两个新的warp投票功能，并且支持GT200微体系结构(Microarchitecture)。



版本1.3增加了对双精度浮点运算的支持。



目前，市面上的GeForce GTX 280、GTX 260以及Tesla S1070、C1060等产品都是版本1.3规范的，而却没有附和版本1.2规范的产品。这意味着Nvidia将在中低端市场上发布不支持双精度浮点运算的产品以降低成本。


体系架构GPU与CPU的一个最根本区别是：GPU与其存储器子系统（显存）是绑定在一起的。CPU的存储器子系统被设计成了可扩展的—想加内存，只需要往DIMM槽里插一条新的；而GPU的显存则是直接焊在PCB板上，提供了更好的电气特性，从而获得了更高的性能。
GT200是一种面向高端市场的单芯片GPU。曾经的高端GPU G80如今已经成为了面向主流市场的中端产品。随着技术发展，GT200也将逐渐从旗舰型号演变为产品线中较低等级的型号。而ATI的策略则是设计一款面向主流市场的GPU—RV770，再用双GPU显卡来获得更高的性能。
两种方案各有千秋。使用单芯片可以获得更高的性能，但在良品率和单位成本上没有优势。此外，太大的晶圆面积也不能像小尺寸晶圆芯片一样可以通过单卡双核的方式面向不同的市场定位。然而要断言单片方案和双片方案的优劣则为时尚早，两者的表现在不同的应用场合各有高下。
但在通用计算领域，单片方案则毫无疑问具有更大优势。单片GPU比两块封装在一起的小GPU更有用处。对CPU来说，多核处理对编程模型几乎没有什么影响。CPU的cache本身就有数据一致性（Coherency & Consistency ,简称CC，即CPU内cache的数据与CPU外部任何存储单元的数据保持一致），因此在几个CPU之间实现数据一致性也没什么难度；而GPU则避开了Cache数据一致性问题的开销。既然在一个GPU内也没有一致性，多个GPU之间的一致性就更加无从谈起了。多个GPU之间无法通信，除非开发者自己设计一套体系来管理数据共享和通信。然而即使使用像CUDA这样优秀的计算模型，也很难实现这一功能。Nvidia将GPU用于计算领域的决心，是设计单芯片GPU的出发点，也切实体现了GPU通用计算对Nvidia的重要性。
图三表示了三种高吞吐量处理器的体系架构：G80，G200和Niagara II。值得指出的是两种GPU使用的缓存是只读的纹理缓存，而Niagara II的缓存则是满足一致性要求的。GT200的显存位宽达到了512bit，由8个64bit GDDR3存储器控制器组成，而前一代的显存位宽只有384bit。显存带宽在不同工作模式下有所不同，在存储器控制器和显存工作在1107MHz时，可以提供高达141.7GB/s的峰值带宽，较上代提高了65%。GDDR3存储器控制器在提供高带宽和高位宽的同时，也提供了多种存储器访问模式，因此不仅峰值性能很高，实际使用中的效率也很好。面向高性能消费领域（也就是游戏）的产品拥有1G的显存，而面向专业应用领域的产品则将显存扩展到了4G,但同时也降低了带宽（显存和存储器控制单元工作在800MHz）。大多数的GPU都没有提供ECC（Erro Control Code）校验功能，因此不能用于需要高可靠性的大规模GPU集群中。


GT200与主机（host）的外部接口也的到了升级。PCI-E 2.0 x16插槽提供了为上下行数据各提供了高达8GB/s的带宽。考虑到PCI-E数据封包的影响，实际可用的带宽大约在5-6GB/s。



GT200概览



作为一种计算设备，GT200是一种多核芯片，设计重点是能够在处理数据并行问题时获得很高的数据吞吐量（compute throuput）。为了能够提高芯片整体的性能，单个线程的性能和执行时间做出了牺牲。这与通用处理器形成了反差：通用处理器的设计重点是提高单个线程的性能，减少执行和通信延迟，而不在计算吞吐量上。



GT200的结构可以被分为两层：第一层是10个TPC（Thread Processing Cluster， 线程处理器群），每个TPC又有三个SM（Streaming Multiprocessors,SM，又称Thread Processor Array， TPA）和一个纹理流水线组成。纹理流水线为所在TPC内的三个SM提供存储器流水线。每个SM大致相当于一个具有8路SIMD功能的现代处理器，拥有独立的完整前端，包括取指，译码，发射和执行单元等，但存储器流水线是与同一个TPC中的其他SM共享的。下图中是GT200与其前代产品G80的框图。两代产品的主要不同是G80拥有8个TPC，每个TPC中有两个SM；而GT200的TPC数量增加到了10个，每个TPC中的SM数量也增加到了3个。



ATI和Nvidia的市场部门曾经将GPU中的执行单元（ALU，FPU）称为“核”，并声称GPU拥有数百个“核”,而现代的CPU则只有几个，这直接导致了术语和概念上的混乱。事实上，把SM与CPU的核相比更加合适。和当代的CPU核一样，SM也拥有独立的取指和调度单元构成的完整前端。



GPU通过向量机技术增强了计算性能，减少了控制方面的开销。例如，GT200和G80种的每个SM又包含8个线程处理器（Thread Processor，TP）。线程处理器也有其他的名称，通用计算程序员可能将他们称为流处理器（Streaming Processor）,而图形学程序员也许会将他们称为渲染核（Shader Core）。SP并不是独立的处理器核，它们有独立的寄存器和指令指针，但没有取指和调度单元来构成完整的前端。因此，SP更加类似于当代的多线程CPU中的一条流水线。



计算架构



GPU的任务分配单元（work distribution unit,或者global block scheduler）将线程块一级的粗粒度并行分配到整个芯片上。启动CUDA 核时，需要将grid的信息从CPU传输到GPU。任务分配单元根据这些信息将block分配到SM上。任务分配单元使用的是轮询调度：轮流查看SM是否还有足够的资源来执行新的block，如果有则给SM分配一个新的block，如果没有则查看下一个SM。决定能否分配成功的因素包括：每个block中的thread数量，占用的共享存储器大小，每个thread占用的寄存器大小，以及每个SM剩余的资源。任务分配单元的目标是将线程均匀的分配到每个SM上，从而获得更高的并行执行机会。



SM 概览



GT200的惊人计算能力来自于处理器阵列。处理器被分为两层，而流处理器处于两层结构中较低的一层。SM的框图如下图5所示：



抛开商业上天花乱坠的宣传不说，SM实际上是一种高度线程化的单发射（single-issue） SIMD处理器。和当代的CPU核一样，SM也拥有独立的取指和调度单元构成的完整前端。大多数情况下，控制开销被8个功能单元分摊，但也有一部分的控制逻辑是直接控制单个线程的，无法共享。



每个最多能够同时执行8个block或1024个thread。虽然在CUDA中可以对thread和block进行操作，但要管理SM中的thread仍然十分烦琐，效率也很低。这是因为每个block最多可以有512个thread，粒度太粗，而每个thread的粒度又太细。



任务分配单元可以为每个SM分配最多8个block。而SM中的线程调度单元又将分配到的block进行细分，将其中的线程打成较小的包，称为线程捆（warp）。在CUDA中，warp对程序员来说是透明的，它的大小可能会随着硬件的发展发生变化，在当前版本的CUDA中，每个warp是由32个线程组成的。Warp的大小对操作延迟和访存延迟会产生影响。新一代的SM能够同时处理32个warp，而上一代G80架构的SM则只能处理24个warp,或者768个thread。注意到，两代SM都能够同时处理8个block。



SIMT 编程模型



与现代的微处理器不同，Nvidia的SM没有预测执行机制-没有分支预测单元(Branch Predicator)。GPU的任务是处理不需要预测的任务，因此也就可以省掉预测相关的电路，降低功耗。在需要分支时，只有当warp中所有的线程都计算出各自的分支的地址，并且完成取指以后，warp才能继续往下执行。上面介绍过，一个warp内的线程是由同一个SM执行的，而一个SM只有一组取指、译码单元，这无疑大大的降低了分支的效率。



Nvidia将这种执行模型称为SIMT（Single Instruction, Multiple Thread，单指令多线程），SIMT是对SIMD（Single Instruction, Multiple Data，单指令多数据）的一种变形。两者的区别在于：SIMD的向量宽度是显式的，固定的，数据必须打包成向量才能进行处理；而SIMT中，执行宽度则完全由硬件自动处理了。以CPU中的SIMD机制为例，在SSE指令集中，有专门的指令能够一次处理4个单精度浮点数。那么如果要用这些指令处理一个由单精度浮点数构成的数组，则一定必须要把这个数组中的数据按照四个一组的形式交给CPU来处理。而在CUDA中，可以将这个数组划分为若干个线程块，不论一个线程块里的线程数量是多少（只要不超过512个线程的上限），硬件都能适应。如果按照SIMD模型来设计CUDA，那么我们就必须自己把数据组织成warp的形式，这无疑将大大的降低编程的灵活性。在SIMT中，如果需要控制单个线程的行为，必须使用分支，这会大大的降低效率。例如，如果一个warp中需要进行分支（即warp内的线程执行的指令指针指向不同的位置），性能将急剧的降低。如果一个warp内需要执行N个分支，那么SM就需要把每一个分支的指令发射到每一个SP上，再由SP根据线程的逻辑决定需不需要执行。这是一个串行过程，因此SIMT完成分支的时间是多个分支时间之和。



另外一个重要不同是SIMD中的向量中的元素相互之间可以自由的相互通信，因为它们存在于相同的地址空间（例如，都在CPU的同一寄存器中）；而SIMT中的warp中的每个线程的寄存器都是私有的，它们只能通过shared memory来进行通信。



新的CUDA版本中新增加的一些指令并不符合SIMT的设计思想。前面提到过，SIMT中是由硬件来适应线程块里的线程数量的，那么warp对于编程人员来说应该是透明的。但新增加的warp voting函数中的_any()函数和_all函数的作用范围都只是在一个warp内，而不是一个block内。_any()函数的功能是：如果一个warp中有至少一个线程的执行结果返回真值，则any()函数返回真，否则为假；all()函数的功能是如果一个warp中的所有线程执行结果都为真，则返回真，否则为假。从1.2以上版本的硬件都支持这两条函数，这意味着warp首次以显式的形式出现了。某种意义上来说，这是向传统的SIMD模型的一种倒退。


取指


为了能够详细说明SM的内容,我们对上节的图五重新进行了注释:


The front-end of the SM has a conventional set-associative instruction cache that automatically fills instruction cache lines when a branch target address misses. Currently, NVIDIA has not released any details of precisely where this structure resides, or its organization. The cache management policy, the capacity, access latency and almost all other information is unknown.
SM的前端拥有一个传统的组相连指令缓存(set-associative instruction cache), 在发生分支目标地址不命中(Cache Miss)时, 它能够自动的填满指令缓存线。 Nvidia没有提供关于指令Cache的任何细节。
Warp instructions are fetched into a multithreaded instruction buffer, which probably contains 32 or 64 entries – one to two entries per warp in-flight.
取出的warp指令被放置在一个多线程指令缓冲上。指令缓冲拥有32或者64个入口，每个活动warp能够获得一到两个入口。

发射
The instruction issue logic is responsible for selecting a warp instruction to issue each cycle. The instruction buffer is located in close proximity to the issue logic for the SM; or the issue logic may simply be implemented as additional metadata for each warp instruction in the buffer.

指令发射逻辑电路在每一个时钟周期选择一条warp指令进行发射(Issue)。在每个SM中，指令缓冲在物理上的位置与发射逻辑电路是紧紧连在一起的。
Instructions are scoreboarded to prevent various hazards from stalling execution. When all the operands for an instruction and the destination registers/shared memory are available, the instruction status changes to ‘ready’. Each cycle the issue logic selects and forwards the highest priority ‘ready to execute’ warp instruction from the buffer. Prioritization is determined with a round-robin algorithm between the 32 warps that also accounts for warp type, instruction type and other factors.

为了避免各种可能影响执行的问题，发射逻辑对指令设置了优先级。当一条指令需要用到的寄存器和shared memory资源都处于可用状态时，这条指令的状态将被设置成“准备”。在每个时钟周期，发射逻辑从缓冲中选区中优先级最高的处于“准备”状态的指令。发射逻辑电路在32个warp间使用轮询算法决定指令的优先级。优先级同时也要受warp类型，指令类型和其他一些因素的影响。

A warp which has multiple ready instructions can continue to issue until the scoreboarding blocks further progress or another warp is selected for issue. This means that the scoreboarding actually enables very simple out-of-order completion. A warp could issue a long latency memory instruction, followed by a computational instruction and in that case, the computation would end up writing back its results before the memory instruction. This is a very limited form of out-of-order execution, comparable to techniques used in Itanium and much less aggressive (and more power efficient) than a fully renamed and out-of-order issue processor such as the Core 2.
如果一个warp有多条处于准备状态的指令等待执行，这些指令将被连续发射，直到重新计算状态和优先级，或者发射逻辑选择了另外一个warp进行发射。这意味着优先级策略实现了简单的乱序执行功能。一个warp的指令中可能包括一次延迟很长的访存，然后是一次计算；但在实际运行中，可能计算在访存完成之间就已经结束了。这种乱序执行与现代CPU使用的乱序执行技术相比十分简单，作用也有限，但却是能够极大的节省晶体管和能耗方面的开销。

Registers
The register files are primarily designed with high bandwidth in mind. The GT200 register files are 64KB, compared to 32KB in the earlier G80 generation. The SM has a total of 16K register file entries partitioned across the SPs. Each of the SPs has a 2K entry register file that can be used by up to 128 threads and probably supports 16 or 24 different banks. Register file entries are 32 bits, and the new double precision data types (both integer and floating point) consume two adjacent registers. The register file is dynamically partitioned between thread blocks by the JIT/driver, and within the allocation for each thread block, the registers are statically assigned to a given thread.
An individual thread can have 4-128 registers.

寄存器（register）
寄存器拥有很高的带宽。GT200拥有64KB的寄存器(Register Files)，而G80的寄存器是32K。每个SM中有16K寄存器是能够被SP访问的，每个SP能够平分到2K的寄存器单元。在每个SP上最多可以保存128个线程的信息。。每个寄存器单元(Register Files Entry)的宽度为32bit，最新加入的64bit数据类型（双精度浮点和64位整数型）将占用两个相邻的寄存器单元。JIT/驱动能够动态的为线程块分配寄存器，而每个线程块中的线程占用的寄存器大小则是静态分配的，在线程块寿命期间都不会更改。每个独立线程能够拥有4到128个寄存器单元。 

 

 

共享存储器（Shared Memory）
GT200的每个SM拥有16KB shared memory，用于同一个block内的thread间通信。为了使一个warp内的thread能够并行的访问shared memory， 4096个入口被划分为16个bank，每个bank拥有32bit的宽度。在不发生bank conflict时，访问shared的速度很快，访问延迟与register相同。在不同的block之间，shared memory是动态分配的。在同一个block内，所有的线程都能够访问shared memory。，在CUDA编程模型中，Shared memory是进行线程间低延迟数据通信的唯一渠道，因此其地位至关重要。

1.2版本以后的硬件的一个新特性是能够在shared memory内进行高速的原子操作。这里的原子操作是指保证在每个线程能够独占的访问存储器，即只有当一个线程完成对存储器的某个位置的操作以后，其他线程才能访问这一位置。1.1版本的硬件只能支持对global memory的原子操作。访问global memory需要很长的访存延迟（长达数百个时钟周期），性能很低。在GT200及以后的GPU上，可以支持对shared memory中32bit操作数的原子操作指令（其中包括CAS指令）。对一个寄存器的原子操作很容易实现，以后的GPU也许能够在shared memory里实现对横跨两个寄存器的64bit字进行原子操作。

The majority of loads from memory write to the thread private register files, rather than the shared memory. To simplify the design of the SM, the architects decided not to directly load data from the global memory into the shared memory. Instead, to move values into the shared memory, data must be loaded into the registers and then separately moved into shared memory. CUDA programs also use the __syncthreads() function to synchronize threads of a thread block at
inter-thread communication points. The SM implements this with a single barrier synchronization instruction that synchronizes up to 512 concurrent threads in the thread block with low latency.

在实际使用中，从显存中读出的数据被写到register中而不是shared memory中。因此，为了简化设计，GPU的结构里shared memory和global memory没有直接连接起来。要把global memroy中的数据写到shared memory中，必须先把数据写到寄存器里，再转移到shared memory中。


执行单元

当代GPU必须拥有丰富的执行资源和强大的计算能力。Nvidia公司通过SIMT-一种SIMD的变形实现了这一目标。如上文所述，SIMT在提供了与SIMD相同的性能的同时，还隐藏了复杂的架构；如果能够使warp内所有线程的执行的分支相同，SIMT执行条件分支语句的性能也比SIMD要好一些。



图五中的高速计算单元集群的运行频率是取指单元、调度单元、寄存器或共享存储器的两倍。在每个时钟周期（对高速计算单元来说是两个快时钟周期）可以执行一条新的warp指令。



高速计算但能源集群中，最主要的执行资源是8个32bit ALU和MAD(multiply-add units，乘加器)。它们能够对符合IEEE标准的单精度浮点数(对应float型)和32-bit整数（对应int型，或者unsigned int型）进行运算。每次运算需要4个时钟周期（四个快时钟周期，而不是控制逻辑和存储器的核心时钟）。因为使用了四栈流水线，因此在每个时钟周期，ALU或MAD都能取出一个warp 的32个线程中的8个的操作数，在随后的3个时钟周期内进行运算并写回结果。



控制流指令（CMP,比较指令）是由分支单元执行的。如前文所述，GPU没有分支预测，因此在分支得到机会执行之前，它将被挂起。和算术指令一样，一个分支warp指令也需要4个时钟周期来执行。



除了标准的功能单元以外，每个SM还拥有两个能够执行不那么常用的运算的执行单元。第一种是用来处理寄存器中的64位浮点和整型操作数的64bit乘加单元，每个SM中有一个这样的单元。这种双精度FMA单元能够支持标准的IEEE754R对双精度操作数的要求，可以进行异常处理，也能够进行64bit整型算术。它能够完成带有舍入的乘加运算，支持高精度的类型转换。由于这样的单元在每个SM中只有一个，因此GPU的双精度计算速度只有单精度速度的1/12-1/8也就不足为奇了。Nvidia已经充分注意到了双精度运算对通用计算的重要性，下一代产品的双精度将会得到很大提高。



第二种是特殊函数单元，用来执行一些特殊指令。SFU用来执行超越函数，插值，倒数，平方根倒数，正弦，余弦以及其他特殊运算。SFU执行的指令大多数有16个时钟周期的延迟，而一些由多个指令构成的复杂运算，如平凡根或者指数运算则需要32甚至更多的时钟周期。SFU中用于插值的部分拥有若干个32-bit浮点乘法单元，可以用来进行独立于FPU的乘法运算。SFU实际上有两个执行单元，每个执行单元为SM中8条流水线中的4条服务。向SFU发射的乘法指令也只需要4个时钟周期。



既然有了这些执行单元，那么当然也需要一套机制来从显存读写数据。读取和存储的指令会从SM控制其发射到纹理流水线。关于纹理流水线的内容，我们将在下文中讨论。



Dual 'Issue' 并发执行



Nvidia的微架构涉及中，吞吐量与延迟之间的关系很有意思。CPU执行一条指令只需要一个周期，而SP执行一条指令再怎么也有至少4个周期的延迟。因为SM每两个“快周期”（执行单元经过倍频，周期是核心周期的一半）就能发射一条指令，那么就有可能让多个指令并发执行。这种能力被Nvidia称为‘dual issue’-虽然实际上是让不同种类的功能单元能够同时运行。SP单元在执行4周期指令的同时，其他的执行单元也能执行其他的warp指令。



如图六所示,SM每两个时钟周期就能发射一条warp指令。在第一个周期，一条MAD指令被发射到FPU单元。两个时钟周期以后，一条MUL指令被发射到了SFU单元。又过了两个时钟周期以后，FPU单元开始执行另一条MAD指令。再过了两个时钟周期，SFU开始执行一条占用很长时间的超越函数指令。应用这项技术能够使渲染核心的计算吞吐量提高50%，同时每两个时钟周期只需要发射一条指令就能满足需要，大大降低了标志设置和优先级逻辑的复杂度。不是所有的指令组合都能够并发执行。例如，双精度浮点单元和单精度浮点单元共享了一部分逻辑，因此无法同时使用。



来自寄存器设计的压力



SM中的8个SP需要访问寄存器，拥有很多bank的shared memory需要访问寄存器，SM中的计算单元不仅能够并发执行，时钟还经过了倍频，这一切给寄存器的设计带来了很大的麻烦。在每个时钟周期，执行单元最多能够并发执行8条FMAD指令和8条FMUL指令。每条FMAD指令需要三个输入操作数，一个输出操作数；每条FMUL指令需要两个输入操作数，一个输出操作数。所以在一个快时钟周期内，最多需要40个输入和16个输出。而且，寄存器单元和shared memory还不是工作在快时钟上，而是更慢的核心时钟上。因此寄存器和shared memory必须能够提供80个输入，32个输出，这意味着整个存储器阵列必须提供112个端口！由于还存在像纹理单元（texturing unit）和存储流水线（load store pipelines），因此实际上的端口总数更加接近128。


纹理，渲染和存储器流水线


现代GPU使用纹理流水线和渲染流水线进行数据的输入输出。CPU上的读取和存储单元是核心中密不可分的一个部分，而GPU中的纹理和渲染输出流水线则与GPU的计算核心是独立的。在GT200中，每个TPC（Thread Prcessing Cluster, 线程处理集群）包括三个SM，一条纹理流水线和一个与渲染输出单元（render output units， ROP）通信的端口。由于本文将GPU作为计算设备来介绍，因此也就不去深究这两个组件在图形学的作用，而把重点放在如何在通用计算中使用它们了。图七是GT200的读取和存储流水线的结构。



读取（Load）和存储（Store）指令是由SM生成的，但当这些指令的发射和执行都在一个工作在完全不同时钟域的硬件上。首先，读取和存储指令被发送到SM控制器。SM控制器负责管理对存储器流水线访问，同时也协调纹理单元和SM的不同时钟。纹理流水线和读取流水线共享了一部分硬件，所以不能同时使用。



读取的第一步是计算地址-地址可以由存储器访问指令中的寄存器中的值与地址偏移量相加得到；然后将计算得到的40b虚拟地址（在G80中是32bit）转换为MMU实际使用的物理地址；最后，向整个warp发射读取指令，并且将指令通过片内互联总线发送到连接着GDDR3存储器控制器。
储存命令的执行与之类似：首先需要计算地址；然后向通过片内互联总线发送到ROP单元，再发送到GDDR3存储器控制器。原子指令同样也需要经过存储流水线和ROP单元发送。虽然每次存储器访问时，指令是按照warp发射的，但存储器控制器是按照half warp执行这些指令的，也就是一次只能够完成16个访问。


The memory pipeline depicted in Figure 7 features two specialized texture caches. These caches are distinct from a traditional CPU cache in several ways. First of all, CPU caches have locality in a single dimension because memory addressing for most architectures is linear. When requesting a data word (which might be 4-8B), an entire 64B cache line will be fetched - so on top of the requested data, there is another 56-60B of data that is brought into the cache because most of the time, this data will be used in close temporal proximity to the original requested data. Textures are fundamentally 2 dimensional objects and are stored in memory in a continuous fashion with respect to both the X and Y dimensions; again traditional data has only one dimension and is only continuous with respect to that dimension. Consequently, texture caches must have 2-dimensional locality to effectively cache textures. Typically, the memory controller is responsible for mapping the 2D texture memory space into one dimensionwith space filling curves before reaching the cache, but the texture caches may have modifications which improve locality and performance.
图七中，我们可以发现存储器流水线拥有两个纹理缓存。这些缓存与CPU的缓存有很大的不同。首先，CPU的缓存往往是一维的，因为大多数的架构中的存储器地址是线性的。当访问一个只有4-8Byte的数据字时，会取出一个缓存单元中所有的64B数据。因为CPU处理的数据往往有很强的相关性，因此多取出的相邻的50-60Byte数据也有可能会被用到。CPU处理的数据是只有一维，因而其缓存也只是在一个维度上是连续的；GPU需要处理的纹理则是连续的二维图像，因此纹理缓存也必须是在两个维度上连续分布的。典型的存储器控制器会将二维的纹理存储器空间映射为意味
Second, texture caches are also read-only and have no coherency. When a texture is written, the entire texture cache hierarchy must be invalidated, rather than tracking the validity of individual data within the address space. Third, texture caches are used to save bandwidth and power only - in contrast to CPU caches which are also essential for lowering latency. Texture caches cannot service requests out-of-order and do not meaningfully impact latency (in comparison a CPU cache might lower latency from 100ns to 7ns). The L1 texture caches in the GT200 reside in the TPC and are 24KB, but they are partitioned into 3x8KB caches in each TPC. The L2 texture caches are located with the memory controllers, each one is 32KB for an aggregate 256KB across the entire device.
其次，纹理缓存是只读的，也不满足数据一致性。当纹理被修改以后，必须更新整个纹理缓存，而不是纹理缓存中被修改的一小部分。
第三，纹理缓存的主要功能是为了节省带宽和功耗，而CPU的缓存则是为了实现较低的延迟。
GT200中，每个TPC拥有的一级纹理缓存是24KB，但是被分为三个8KB的块。L2纹理缓存位于存储器控制单元中，每个大小是32KB，所以整个器件已共有256KB。
因为对global的存储器访问没有缓存，因此显存的性能对GPU至关重要。为了能够高效的访问显存，读取和存储必须对齐，宽度为4Byte。如果没有正确的对齐，读写将被编译器拆分为多次操作，极大的影响效率。此外，多个half-warp的读写操作如果能够满足合并访问（coalesced access），那么多次访存操作会被合并成一次完成，从而提高访问效率。
GT200的一个经常被忽略的重大改进是对存储器的合并访问的约束条件比上一代架构放宽了很多。
在1.0和1.1器件的合并访存条件十分严苛。首先，访存的开始地址必须对齐：16x32bit的合并必须对齐到64Byte（即访存起始地址必须是64Byte的整数倍）；16x64bit的合并访存起始必须对齐到128Byte；16x128bit合并访存的起始地址必须对齐到128Byte，但是必须横跨连续的两个128Byte区域。其次，只有当第K个线程访问的就是第K个数据字时，才能实现合并访问，否则half warp中的16个访存指令就会被发射成16次单独的访存。
在新的1.2版本以上的硬件中，不仅条件放宽了许多，而且还能支持对8bit和16bit数据字的合并访问（分别使用32Byte和64Byte传输）。首先在一次合并传输的数据中，并不要求线程编号和访问的数据字编号相同。其次，当访问128Byte数据时如果地址没有对其到128Byte，在老版本硬件中会产生16次访存指令发射，而在新版本中只会产生两次合并访存。而且，这两次合并访存并不是两次128Byte的。例如，一次128Byte访存中有32Byte在一个区域中，另外一个区域中有96Byte，那么只会产生一次32Byte合并访存(对有32Byte数据的区域)和一次128Byte（对有96Byte数据的区域）。