GPU&CUDA几个基本概念

　　本文介绍 GPU 和 CUDA 相关的几个基本概念：SP、SM、warp、thread、block 和 grid。其中 SP，SM，warp 是硬件（GPU hardware）概念，而 thread，block，grid 是软件（CUDA）概念。

参考了以下几篇博客或文章：
1. 　http://www.myexception.cn/cuda/1929840.html
2. 　http://blog.csdn.net/junparadox/article/details/50540602
3. 　http://www.cnblogs.com/dwdxdy/p/3215158.html

硬件相关

SP （Streaming Processor，流处理器）：最基本的处理单元， 最后具体的指令和任务都是在 SP 上处理的。GPU 进行并行计算，也就是很多个 SP 同时做处理。SP 也可以称为 CUDA Core 或者 thread，一个 SP 对应一个 thread。

SM （Streaming Multiprocessor，流多处理器）：多个 SP 加上其他的一些资源组成一个 SM，GPU 硬件的并行性就是由 SM 决定的。以 Fermi 架构为例，其包含以下主要组成部分：
　 -　CUDA cores
　 -　Shared Memory/L1 Cache
　 -　Register File
　 -　Load/Store Units
　 -　Special Function Unit（SFU）
　 -　Units Warp Scheduler

warp（线程束）：一个 SM 中的所有 SP 在物理上分成 warp 的，共享同一个 memory 和 instruction unit。warp 是 GPU 执行程序时的调度单位，目前 CUDA 的 warp 的大小为 32，这个 warp 中的 32 个 threads(SP) 是一起工作的，执行相同的指令。也就是当进行并行计算时，线程数尽量为 32 的倍数，如果线程数不是 32 的倍数，假如是 1，则 Warp 会生成一个掩码，当一个指令控制器对一个 Warp 单位的线程发送指令时，32 个线程中只有一个线程在真正执行，其他 31 个进程会进入静默状态。

软件相关

thread（线程）：一个 CUDA 的并行程序会被以许多个 threads 映射到 SP 上来执行，一个 thread 映射到一个 SP。

block（线程块）：多个线程会组成一个线程块，同一个块内的线程可以同步，也可以通过 shared memory 通信。由于块内的所有线程必须存在于同一个处理器核心中且共享该核心有限的存储器资源，因此，一个块内的线程数目是有限的。然而，一个内核可被多个同样大小的线程块执行，所以总的线程数等于每个块内的线程数乘以线程块数。

grid（线程网格）：多个线程块被组织成一维、二维或三维的线程网格，一个网格内的线程块数往往由被处理的数据量而不是系统的处理器数决定，前者往往远超后者。

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

软硬件映射关系

软件	硬件
thread	SP/CUDA Core
32 thread	warp
thread block	SM
thread grid	Device/GPU

总结

　　从硬件架构来看，在 GPU 中最小的硬件单元是 SP（CUDA Core/thread），而硬件上一个 SM 中的所有 SP 在物理上是分成了多个 warp，warp 中的 SP 可以同时工作的，执行相同的指令，也就是说取指令单元取一条指令同时发射给 warp 中的所有的 SP（假设 SP 都需要工作）。可见，在硬件上一个 GPU -> SM -> warp -> Core/SP。

　　从软件组织来看，最小的逻辑单位是一个 thread，最小的执行单位是 warp，一个 warp 包含一定数目（比如 Tegra 32个)的 SP（thread），因此最好按照这个数目来组织 thread，否则硬件上该 warp 上有些 SP 是不工作的。可见，在软件上一个 grid -> block（-> warp）-> thread。