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.