取指与发射

　　取指和发射单元在物理上是紧密相连的。取指单元从DRAM中取出指令，并装载于指令缓存中。由于GPU没有分支预测单元，因此在出现分支或者循环时有可能会发生指令缓存不命中。指令缓存的设计属于商业秘密，因此没有更多细节可供参考。

　　处理器频率大约是GPU核心频率的两倍多，因此SM中的8个SP每两个处理器周期才能访问一次存储器，或者接受一条新的指令。也就是说，很多操作是以每16个线程为一组完成的。16在CUDA中是一个神奇的数字，在考虑性能优化时half-warp比warp更加重要。

　　SM中的指令是以warp为单位发射的，在属于同一个warp线程间进行通信不需要进行栅栏同步。虽然SP可以设计成接受一条新的指令然后执行两遍，对应的warp大小是16;但实际上SP接受一条指令执行了四遍，warp大小就成了32。将Warp大小定为32而不是16是综合多方面因素的结果，其中一个原因是为了实现双发射超标量并行，我们将在稍后进行介绍。

　　那么，一个SM中的所有active warp又是按照什么顺序发射的呢?为了避免各种可能影响执行的问题，发射逻辑对指令设置了优先级。当一条指令需要用到的寄存器和shared memory资源都处于可用状态时，这条指令的状态将被设置为就绪态。在每个时钟周期，发射逻辑从缓冲中选区中优先级最高的就绪态指令。发射逻辑电路使用一个加权算法计算各个active warp指令的优先级。优先级要受warp类型，指令类型和其他一些因素的影响。

　　如果一个warp中有多条处于就绪状态的指令等待执行，这些指令将被连续发射，直到重新计算状态和优先级，或者发射逻辑选择了另外一个warp进行发射。这意味着优先级策略实现了简单的乱序执行功能。一个warp的指令中可能先存在一次延迟很长的访存，然后是一次计算;但在实际运行中，如果计算不依赖访存得到的数据，可能出现计算在访存完成之间就已经结束的情况。这种乱序执行与现代CPU使用的乱序执行技术相比非常简单，作用也相对有限，但却是能够极大的节省晶体管和能耗方面的开销。

　　寄存器(register)

　　寄存器拥有很高的带宽。计算能力1.2/1.3硬件的每个SM拥有64KB的寄存器(Register Files)，而计算能力1.0/1.1硬件的每个SM只有32KB寄存器。寄存器的基本单位是宽度为32bit(4Byte)的寄存器单元，这样计算能力1.2/1.3的每个SM就有16K个寄存器单元，每个SP能够平分到2K个寄存器单元。每个寄存器单元的宽度为32bit，所以64bit数据类型(双精度浮点和64位整数型)将占用两个相邻的寄存器单元。JIT/驱动能够动态的为线程块分配寄存器，而每个线程占用的寄存器大小则是静态分配的，在线程块寿命期间都不会更改。这就是说，每个线程占用的寄存器资源是由这个线程使用最多寄存器的时刻决定的。每个独立线程能够拥有4到128个寄存器单元。

　　共享存储器(Shared Memory)

　　GT200的每个SM拥有16KB shared memory，用于同一个block内的thread间通信。为了使一个half-warp内的thread能够在一个周期内并行的访问shared memory， shared memory的4096个入口被组织成为16个bank，每个bank拥有32bit的宽度。如果half-warp内有若干个线程访问的数据处于同一个bank中，就有可能出现bank conflict。当不发生bank conflict时，访问shared的延迟与register相同,否则多个线程就要串行的对处于同一Bank中的数据进行访问。在不同的block之间，shared memory是动态分配的。在同一个block内，所有的线程都能够访问shared memory。CUDA编程模型中的Shared memory是进行线程间低延迟数据通信的唯一方法，因此其地位至关重要。

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

　　在实际使用中，大多数从显存中读出的数据都会被写到register中，而不是写到shared memory中。因此，为了简化设计，GPU的结构里shared memory和global memory没有直接连接起来。要把global memroy中的数据写到shared memory中，必须先把数据写到寄存器里，然后才能转移到shared memory中。编译器可以透明的完成这一过程，使用cuda C编程时可以将global memory中的值直接赋给shared memory。

　　由于各个warp在CUDA中的执行顺序并没有什么规律，因此各warp执行的进度也不尽相同。因此使用shared memory进行线程间通信时，必须进行调用一次__syncthreads()函数进行同步。调用syncthreads()内联函数将进行过一次栅栏同步(barrier)，它保证属于同一个block中的所有warp都完成栅栏同步前的操作后才能继续进行。SM可以完成对512个线程的栅栏同步，因此一个Block中最多只能有512个线程。

　　执行单元

　　当代GPU必须拥有丰富的执行资源和强大的计算能力。Nvidia公司通过SIMT-一种SIMD的变形实现了这一目标。如上文所述，SIMT在提供了与SIMD相同性能的同时，还可以通过改变线程数量适应不同的执行宽度。

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

　　Tesla架构中最主要的执行资源是8个32bit ALU和MAD(multiply-add units，乘加器)。它们能够对符合IEEE标准的单精度浮点数(对应float型)和32-bit整数(对应long int型，或者unsigned long int型)进行运算。ALU和MAD运算至少需要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用来执行超越函数，插值，倒数，平方根倒数，正弦，余弦以及其他特殊运算。CUDA中提供了一些带有__前缀的函数，其中一部分就是由SFU执行的。SFU执行的指令大多数有16个时钟周期的延迟，而一些由多个指令构成的复杂运算，如平方根或者指数运算需要32甚至更多的时钟周期。SFU中用于插值的部分拥有若干个32-bit浮点乘法单元，可以用来进行独立于FPU的乘法运算。SFU实际上有两个执行单元，每个执行单元为SM中8条流水线中的4条服务。向SFU发射的乘法指令也只需要4个时钟周期。