Cuda编程-06 硬件执行补充

物理概念

SP流式单处理器

        最基本的处理单元，Streaming Processor，最后具体的指令和任务都是在SP上处理的。GPU进行并行计算，也就是很多个SP同时做处理。现在SP的术语已经有点弱化了，而是直接使用Thread来代替，一个SP对应一个Thread。流式单处理器SP表达的是物理层面概念，线程Thread表达的是逻辑层面概念。

SM流式多处理器

        多个SP加上其他的一些资源组成一个SM, Streaming Multiprocessor，其他资源也就是存储资源，共享内存，寄储器等。流式多处理器SM对应线程块threadBlock，SM表达物理层概念，块表达逻辑层概念。

Wrap线程束

        Wrap是SM调度和执行的基础概念，同时也是一个硬件概念，注意到Wrapp实际上是一个和硬件相关的概念，通常一个SM中的SP(thread)会分成几个wrapp(也就是SP在SM中是进行分组的，物理上进行的分组)，每一个Wrap中有32个thread。这个Wrap中的32个thread(sp)是一起工作的，执行相同的指令，如果没有这么多thread需要工作，那么这个Wrap中的一些thread(sp)是不工作的。

        每一个线程都有自己的寄存器内存和local memory，一个warp中的线程是同时执行的，也就是当进行并行计算时，线程数尽量为32的倍数，如果线程数不上32的倍数的话；假如是1，则warp会生成一个掩码，当一个指令控制器对一个warp单位的线程发送指令时，32个线程中只有一个线程在真正执行，其他31个线程会进入静默状态。

逻辑概念

        thread–>block–>grid：在利用cuda进行编程时，一个grid分为多个block，而一个block分为多个thread。其中任务划分会影响最后的执行效果。划分的依据是任务特性和GPU本身的硬件特性。GRID,BLOCK,THREAD是软件概念,而非硬件的概念。

        从硬件角度讲，一个GPU由多个SM组成（当然还有其他部分），一个SM包含有多个SP（以及还有寄存器资源，shared memory资源，L1cache，scheduler，SPU，LD/ST单元等等）。SP或者Core目前也称为CUDA CORE，而SM目前也称为MP，在KEPLER架构（SM3.0和3.5）下也称为SMX。

        从软件角度讲，GRID，block，thread是thread的组织形式。最小的逻辑单位是一个thread，最小的硬件执行单位是thread warp（简称warp），若干个thread（典型值是128~512个）组成一个block，block被加载到SM上运行，多个block组成整体的GRID。

        这里为什么要有一个中间的层次block呢？这是因为CUDA通过这个概念，提供了细粒度的通信手段，因为block是加载在SM上运行的，所以可以利用SM提供的shared memory和__syncthreads()功能实现线程同步和通信，这带来了很多好处。而block之间，除了结束kernel之外是无法同步的，一般也不保证运行先后顺序，这是因为CUDA程序要保证在不同规模（不同SM数量）的GPU上都可以运行，必须具备规模的可扩展性，因此block之间不能有依赖。

从上面的表述中可以总结：

        在GPU中最小的硬件单元是SP(这个术语通常使用thread来代替),而硬件上一个SM中的所有SP在物理上是分成了几个WraP(每一个wrap包含一些thread),wrap中的SP是可以同时工作的，但是执行相同的指令，也就是说取指令单元取一条指令同时发射给WARP中的所有的SP(假设SP都需要工作，否则有些是idle的).可见，在硬件上一个SM->WARPS->threads(sp).

这就是CUDA的两级并行结构。

        总而言之，一个kernel对应一个GRID，该GRID又包含若干个block，block内包含若干个thread。GRID跑在GPU上的时候，可能是独占一个GPU的，也可能是多个kernel并发占用一个GPU的（需要fermi及更新的GPU架构支持）。