GPU编程初探

GPU（Graphic Processing Unit），图像处理处理器，俗称显卡，主要处理图像、显示等任务（数据运算）

CPU：逻辑运算

GPGPU（General Purpose computing on Graphic Processing Unit ），通用GPU，主要处理通用计算任务。

GPU性能指标：

1. 核心数
2. GPU显存容量
3. GPU计算峰值
4. 显存带宽

CPU与GPU架构差异

CPU+GPU异构架构

GPU不能单独计算，CPU+GPU组成异构计算架构；

CPU起到控制作用，一般称为主机（Host）；

GPU可以看作CPU的协处理器，一般称为设备（Device）；

主机和设备之间内存访问一般通过PCLe总线链接；

NVIDIA GPGPU发展历史

Tesla 架构

1. 代表性GPU：G80、G200
2. 第一款支持CUDA的GPGPU
3. 第一款使用标量线程处理器的GPU
4. 首次引入了单指令多线程（SIMT）执行模型，多个线程使用一条指令并发执行
5. 首次引入共享内存，可编程Cache，线程间通信

 

Fermi 架构（Compte Capability 2.0,2.1）

1. 代表性GPU：C2050、C2070
2. 第一款真正支持通用计算的GPU
3. 提高双精度性能、ECC支持
4. Ture Cach Hierarchy
5. 共享内存与L1 Cache 可配置
6. 大幅提升双精度性能
7. 我国的天河1A超算，采用了C2050 GPU，全球首台采用异构架构的超算，开启异构计算时代

 

 

 Kepler 架构（Compte Capability 3.0~3.7）

1. 代表性GPU：K20、K40、K80
2. 引入DP Unit，双精度性能首次超过1T Flops
3. 减少SM（SMX）数量
4. 大幅增加每个SMX中的CUDA Core数量
5. 192个CUDA Core+64个DP+32个SFU
6. 美国橡树岭国家实验室超级计算系统Titan采用了Kepler架构的GPU
7. 进军科学计算的雄心 

 Maxwell 架构（Compte Capability 5.0~5.3）

1. 代表性GPU：Tegra TX1
2. 去除DP Unit，追求功效比
3. 将SMM内的CUDA Core进行分组，每组有专用的指令调度单元
4. SMM内CUDA Core数量减少为128个，方便调度和控制
5. Tegra TX1 终端高性能ARM+GPU异构计算平台
6. 尝试进入自动驾驶领域

 

 Pascal 架构（Compte Capability 6.0~6.2）

1. 代表性GPU：P40、P100
2. DP Unit 又回来了，重视双精度
3. SM内进一步精简，但每个GPU包含的SM多了
4. NVLink横空出世
5. Global Memory 由GDDR5替换成了HBM2， 带宽大幅增加
6. 面向深度学习的初步定制：CUDNN
7. 硬件直接支持FP16

 

 Volta 架构（Compte Capability 7.0~7.2）

1. 代表性GPU：V100
2. Tensor Core 来了，深度学习定制
3. FP32和INT32两组运算单元相互独立,可同时执行一条FP32指令和一条INT32指令
4. 每个SM单元，4组Process Block
5. 硬件直接支持FP16
6. NVIDIA DGX系列AI超级计算机横空出世

 

 

Turing 架构（Compte Capability 7.5）

1. 代表性GPU：T4
2. Volta架构的小加强
3. 增加了Int8/Int4的推理能力
4. 去除了FP64，深度学习方面专注推理
5. LD/ST砍半
6. 提升管线渲染
7. 增加TR Core，增强光线追踪能力

 

 Ampere 架构(Compte Capability 8.0）

1. 代表性GPU：A100
2. 主要改进是Tensor Core，在深度学习的道路上越走越深
3. 增强了Tensor Core中混合精度FMA指令的计算能力，由每个时钟周期执行64个FMA提升到每个时钟周期执行256个FMA
4. Tensor可支持的数据类型扩展到了FP16、BF16、TF32、FP64、INT8、INT4等数据类型
5. 增加结构化稀疏运算能力

 

GPU硬件抽象

大规模细粒度并行处理器(众核)

层次化的硬件架构：

GPU -> SM(SMX、SMM）->PE(CUDA Core, FP32/FP64）

层次化的内存组织形式：

全局内存(Global Memory)->共享内存(Shared Memory)->私有内存（寄存器）

线程组织

层次化的线程组织形式：

Grid->Block->Thread

线程索引(dim3)

Local_id: (threadIdx.x, theadIdx.y)

Block_id: (blockIdx.x, blockIdx.y)

Block_dim:(blockDim.x, blockDim.y)

Global_id:

X: blockIdx.x * blockDim.x + threadIdx.x

Y：blockidx.y * blcokDim.y + theadIdx.y

线程开启数量和线程组织方式由用户决定

线程与硬件的映射

• 每个Thread都有自己的私有内存（寄存器）
• 同属于一个Block内的所有Thread可通过共享内存（sharead memory）进行通信
• 同属于一个Grid内的所有Thread可通过全局内存进行通信

异构编程

异构计算系统：

• Host（主机端）: CPU
• Device（设备端）: GPU

异构计算程序：

• 始于Host，终于Host：Device程序由Host端开启，Device完成计算任务后，将计算结果传回Host端。
• Host和Devie可以并行执行不同计算程序

 

 CPU与GPU之间的数据传输

1. 独立GPU：CPU内存和GPU内存相互独立且分离
2. Host端除执行一般C语言程序的内存分配外，还需执行如下操作：

• 分配设备内存：cudaMalloc
• 执行Host与Device间的数据传输： cudaMemcpy

 

CUDA程序组成

Host端程序（管理者）

• 运行在CPU
• 内存管理、Host与Device间通信、任务管理

Device端程序（工人）

• 运行在GPU
• 计算任务

编译：NVCC编译器