CUDA简介

CUDA：Compute Unified Device Architecture，统一计算设备架构，CUDA™是一种由NVIDIA推出的通用并行计算架构，该架构使GPU (Graphics processing unit) 能够解决复杂的计算问题。它包含了CUDA指令集架构（ISA）以及GPU内部的并行计算引擎。

CUDA 即是 NVIDIA 的 GPGPU 模型，它使用C语言为基础

和CPU的比较

优点：

• 更大的内存带宽
• 更多的执行单元，虽然频率比CPU低
• 价格更低

缺点：

1. 运算单元多，只适合高度并行化的工作
2. 对于具有高度分支的程序，效率会比较差

The GPU Devotes More Transistors to Data Processing

具体来说，GPU特别适合大量并行的数据运算（高运算密度）。由于对每个数据进行相同的操作，所以对复杂的流控制需求较低，并且因为处理许多数据单元并且具有高运算密度，内存读取延时可以通过运算来隐藏，CPU采用的是高速缓存cache来缩减延时 （latency ）

可伸缩Scaleable的编程模式

核心是三个关键的抽象：

1. 线程组的层次结构
2. 共享内存
3. 障碍同步

这种可扩展的编程模型允许市场上的各种GPU架构，GeForce，Quadro，Tesla等。

Automatic Scalability

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程序分配线程不足进行处理

Grid of Thread Blocks

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.

支持CUDA的GPU列表
计算能力的技术规格