CUDA：Compute Unified Device Architecture，统一计算设备架构，CUDA™是一种由NVIDIA推出的通用并行计算架构，该架构使GPU (Graphics processing unit) 能够解决复杂的计算问题。它包含了CUDA指令集架构（ISA）以及GPU内部的并行计算引擎。
CUDA 即是 NVIDIA 的 GPGPU 模型，它使用C语言为基础
和CPU的比较
优点：
- 更大的内存带宽
- 更多的执行单元，虽然频率比CPU低
- 价格更低
缺点：
- 运算单元多，只适合高度并行化的工作
- 对于具有高度分支的程序，效率会比较差
具体来说，GPU特别适合大量并行的数据运算（高运算密度）。由于对每个数据进行相同的操作，所以对复杂的流控制需求较低，并且因为处理许多数据单元并且具有高运算密度，内存读取延时可以通过运算来隐藏，CPU采用的是高速缓存cache来缩减延时 （latency ）
可伸缩Scaleable的编程模式
核心是三个关键的抽象：
- 线程组的层次结构
- 共享内存
- 障碍同步
这种可扩展的编程模型允许市场上的各种GPU架构，GeForce，Quadro，Tesla等。
A GPU is built around a scalable array of multithreaded Streaming Multiprocessors (SMs).
Streaming Multiprocessors (SMs)
在一个multiprocess中，一个线程块的所有线程可以同时执行，多个线程块block也可以同时执行，使用SIMT (Single-Instruction, Multiple-Thread) 构架，一个线程内部通过指令流水进行指令级别的并行，通过硬件多线程进行线程级别的并行。
一个GPU中有多个SM,每个SM有多个core（processor），但是只有一个指令单元，同时只能够执行完全相同的指令集。
Warp
Multiprocessor执行线程块时，把他们以32个并行线程为一组，称为warps,每个warp由warp scheduler调度执行，warp中的每个线程在相同的程序地址处开始执行，但他们有自己的指令地址计数器和寄存器状态。
当他们执行相同的指令时可以达到最大的效率，但当因为由于依靠数据决定的条件分支产生了分歧（warp divergence），warp连续执行每个分支路径，禁止不相关的线程，执行完成后，又回归到相同的执行路径。不同warps执行互相独立。
多线程的CUDA程序具有自动适应性，它被分解为相互独立的线程块，可以以任意的顺序执行，不管是并行还是串行，所以一个编译好的CUDA程序可以在任意数目的多处理器上运行，只有运行的系统需要知道实际的处理器数量
Kernels
一种扩展的C函数，调用时通过N个不同的CUDA线程执行N次，不像通常C函数一样只执行一次
__global__说明符来声明kernel函数，执行的次数通过<<<...>>>执行配置语法来说明
每个线程对应的唯一标识ID通过内置的threadIdx变量获取
// Kernel definition
__global__ void VecAdd(float* A, float* B, float* C)
{
int i = threadIdx.x;
C[i] = A[i] + B[i];
}
int main()
{
...
// Kernel invocation with N threads
VecAdd<<<1, N>>>(A, B, C);
...
}
Thread Hierarchy
线程被分成成块的网格。Grid和Block都是三维的
例如：
dim3 dimBlock(8, 8, 8);
dim3 dimGrid(100, 100, 1);
Kernel<<<dimGrid, dimBlock>>>(…);
不过通常我们使用的一维grid和block
Kernel<<<block_count, block_size>>>(…);
每个block中的thread数目有限，因为他们在同一个内存核心中，共享有限的内存资源，通常最多1024个线程。
grid中的block的数量：最大数目一般为65535，由被处理的数据的大小或者系统中处理器的数量决定
- 如果超出了任何一个数目，GPU会出错或者产生垃圾数据
- GPU编程的部分时间就是用来处理硬件限制的
- 这个限制意味着必须要对Kernel程序分配线程不足进行处理
index of thread 和 thread ID的关系
threadIdx是一个三维矢量，所以线程可以通过一、二、三维线程索引index来识别，组成对应维度的线程块thread block。
一维索引直接对应线程ID
二维(Dx, Dy)的块， 线程索引为 (x, y)对应ID为(x + y Dx);
三维(Dx, Dy, Dz)块, 索引为(x, y, z) 对应ID为 (x + y Dx+ z Dx Dy).
blockIdx 索引变量可以用来标识网格中的块，blockDim可以获取维度
同步 synchronization
线程通信和资源共享：
通过同步来相互协调，调用固有函数__syncthreads()
只在block层次产生作用，类似于C/C++中的 barrier()函数
原子操作 Atomic Operation
执行读-修改-写的原子操作，在全局或共享内存空间中
串行操作
atomic<op>(float *address, float val);
op的范围: Add, Sub, Exch, Min, Max, Inc, Dec, And, Or, Xor
atomicCAS(int *address, int compare, int val)
warp shuffle
CC >= 3.0 在一个warp内的线程交换变量
int __shfl(int var, int srcLane, int width=warpSize);
计算能力 CC
设备的计算能力由版本号代表，也叫做 SM Version。
这个版本号标识了GPU硬件的支持特性，在应用运行时使用，决定在当前GPU上可以实现的硬件特性和指令。它由主版本号X和小版本号Y组成，表示为X.Y
主版本号表示核心构架相同，副版本号表示核心构架的改进，可能增加了新的特性
| The major revision number | Core architecture |
|---|---|
| 5 | Maxwell architecture |
| 3 | Kepler architecture |
| 2 | Fermi architecture |
| 1 | Tesla architecture. |
The Tesla architecture is no longer supported starting with CUDA 7.0.
