一、相关名词
硬件相关：
SP(Streaming Processor,流处理器)：最基本的处理单元（等价于一个ALU），又叫做CUDA Core
SM(Streaming Multiprocessor,流多处理器)：又叫GPU Core，是多个SP和其他一些硬件资源组成的，下面以GTX 580为例讲述SM
一个GTX 580 的SM包含32个SP（CUDA core），2个取址译码单元（Fetch/Decode），和Shared Memory和Execution contexts组成。一个SM上驻扎block个数，warp个数，thread个数上限都是由硬件决定的，另外一个限制条件就是register和shared memory等资源的是否充足，不充足的情况下，驻扎个数达不到上限。
另外，这里SP的个数32与warp的尺寸相等只是巧合，在低版本架构中SP个数可能为8（一个warp要4次才能执行一遍），高版本Kepler架构的GTX 680中，一个SMX中有192个SP，还有48个SP的情况（这种情况下如何执行未深入了解，可能会让我们疑惑，根据规则：一个SM同一时刻只有一个warp在运行，且warp的尺寸为32，，SP小于等于32个我们都能理解，SP没有一个处于限制状态甚至很紧张，那么多于32个SP的SM中，其他SP难道都在闲置中？？？）。参考：http://15418.courses.cs.cmu.edu/spring2016/lecture/gpuarch/slide_056，上述规则可能并不适用，一个SM每个clock选择的warp数，跟Fetch/Decode单元的个数有关，上图中明明白白写着Two warps are selected each clock(decode, fetch and execute two warps in parallel)，SP总共32个，可能是两个warp轮换使用。
warp：warp是GPU执行程序时的调度单位，目前CUDA的warp大小为32，warp是由连续的32个thread组成的，warp可以理解成一种硬件机制，这种机制保证这32个连续thread能够天然同步执行（无时无刻不在执行相同的代码，无需软件同步，遇到divergent 则等待），warp之间切换时零开销的，因为每一个thread的相关数据，执行环境都有自己的专属空间（比如register file， execution contexts），切换warp只需要让sp关联到另一块thread的专属空间即可，没有重新加载的过程。另外，同一时刻一个SM上只能有一个warp在工作。
软件相关：
thread（线程）：每一个thread都有自己专属的register，执行上下文（还不确定，可能有误），local memory等硬件资源，每一个thread都对应一个SP，但是每一个sp则对应多个thread，thread是一个软件概念，它的硬件资源除了SP不是专属的，其他基本都是专属的，一个SM驻扎成千上万个thread，并不是说真的有成千上万个SP(也可能真的有这么多，但绝对不会有thread的数量多)，而是我们将寄存器register，local memory，执行上下文等硬件资源平均分配成 成千上万份，SP只需要在这些不同的thread所对应的专属资源间以warp为单位来回切换即可。
block（线程块）：多个thread组成一个block，block的维度和大小可以通过调用kernel函数时来指定，一个block中的所有thread必须驻扎在同一个SM上，且可以通过软件同步（通过调用__syncthreads()），但不同block之间的thread无法同步，并且核函数中开辟一块shared memory是一个block中所有thread共享的。
grid：多个block组成一个grid，尺寸在调用kernel函数时指定。
二、函数的声明
| Executed on the: | Only callable from the: |
__global__ void KernelFunc() | device | host |
__device__ float DeviceFunc() | device | device |
__host__ float HostFunc() | host | host |
__global__：返回值类型必须是void.
__global__和device函数：
尽量不用递归（不鼓励）
不要使用静态变量
少用malloc（现在允许但不鼓励）
小心通过指针实现的函数调用
三、内存模型
变量声明 | 存储器 | 作用域 | 生命期 |
必须是单独的自动变量而不能是数组 | register | thread | kernel |
自动变量数组 | local | thread | kernel |
__shared__ int sharedVar; | shared | block | kernel |
__device__ int globalVar; | global | grid | application |
__constant__ int constantVar; | constant | grid | application |
四、 SIMD和SIMT的疑问
这里我们以这样一个只包含8个SP的SM为例，说明自己的看法，为仔细认证，可能有误：
上图中上图的SM可能驻扎了许多warp，但执行每一个warp时都需要轮换4次，因为warp有32个thread，而SM只有8个SP；那么在执行一个warp内的32个thread时，可以称之为SIMD，单指令多数据，但是一个GPU肯定有多个SM，有些芯片允许不同的SM允许执行的指令流不同（周斌老师说这是符合芯片的设计哲学的），但是在GPU中不同的SM上执行相同的指令流，可以称之为SIMT，单指令多线程。自己的理解，可能错的离谱。
参考：http://hustcat.github.io/gpu-architecture/
