云栈社区»论坛 › 技术文档「 Note & Doc 」 › Tesla到Pascal架构演进：五代NVIDIA GPU核心特性对比与CUDA能力 ...

发回帖发新帖

1426 积分	0 好友	208 主题

发消息

Tesla到Pascal架构演进：五代NVIDIA GPU核心特性对比与CUDA能力解析

发表于 5 天前 | 查看: 20| 回复: 0

本文将承接上篇，从架构特点、对应产品及CUDA算子支持的变化，详细解析NVIDIA的Tesla、Fermi、Kepler、Maxwell和Pascal共五代GPU架构。

1 Tesla 架构分析

Tesla架构主要包含G80和GT200两个版本，是NVIDIA第一代“统一着色与计算架构”(unified shader and compute architecture)，其架构如下：

其核心内容在上一篇中已有分析。Tesla架构基本奠定了现代GPU运算的雏形，建议延伸阅读关于 AI经典GPU架构---Tesla 的详细介绍。

1.2 Tesla架构对应的产品

2 Fermi 架构分析

Fermi架构于2010年4月发布，采用40nm/28nm工艺，芯片代码“GFXXX”。这代架构的主要亮点包括：

CUDA core的精度运算能力提升；
支持同一个上下文内的kernel并行；
增加ECC功能；
指令分发采用双warp调度；
GigaThread能力增强，上下文切换速度大幅提升；
其它改进：更大的shared memory与L1 Cache、更多的CUDA core、更快的原子操作。

2.1 架构特点

Fermi的主架构图如下所示，逻辑功能主要包括存储、调度、运算和接口四大模块。

从架构图上看，Fermi并未明确绘制TPC（Texture Processing Clusters），但在图形卡设计中依然存在TPC及GPC（Graphics Processing Clusters）。与上一代Tesla GT200相比，Fermi的主要参数优化点对比如下：

核心优化在于引入了支持更高精度的FMA运算、增加了ECC功能、将LD/ST的地址提升到64bit，并支持kernel并发执行。其余提升多为数量与容量的扩展。

2.1.1 SM特点
Fermi架构共包含512个CUDA core，分布于16个SM（Streaming Multiprocessor）中，每个SM包含32个core。此外，每个SM还包含16个LD/ST单元、4个SFU以及两组warp调度与分发单元。

值得注意的是，SM中存在一个Uniform Cache，公开资料未详述其用途，可暂理解为L2 cache的一部分。此外，以计算为主的SM存在两种公开版本，主要区别在于是否包含纹理单元。

2.1.2 CUDA Core
Fermi架构明确了“CUDA Core”的称谓，其最大提升在于FMA（Fused Multiply-Add）运算能力。FMA遵循IEEE 754-2008浮点标准，相比之前的MAD（Multiply-Add）在乘法后保留了完整精度，因此计算结果更精确。

每个CUDA Core在每个时钟周期可完成一次单精度FMA，整卡吞吐量为512 FMA ops/clock；双精度吞吐量为256 FMA ops/clock。相比Tesla GT200的30 FMA ops/clock，性能提升超过8倍。

2.1.3 任务下发与调度
Fermi的任务调度分为两层：

全局调度：由第二代GigaThread引擎负责，其上下文切换速度比Tesla快10倍，并支持同一context内的多个kernel并发执行，提高了计算单元利用率。
SM内部调度：采用Dual Warp Scheduler，每个调度器可独立获取、分发指令。结合运算单元分组（两组CUDA Core、一组SFU、一组LD/ST），支持Dual issue操作，即一个周期内可向两个不同的执行组下发指令，但双精度指令不能与其它指令同时下发。

2.1.4 ECC功能
Fermi是首个在GPU中引入ECC（Error Correcting Code）功能的架构，对寄存器、共享内存、L1/L2缓存均提供保护，可检测并纠正单比特错误，检测多比特错误。ECC虽会略微降低内存效率，但极大地提升了数据可靠性，成为后续Tesla系列产品的标配。

2.1.5 内存结构改进
硬件上，Fermi将L1 Cache与Shared Memory整合在统一的64KB物理块中，比例可配置（如48KB+16KB）。同时，统一了Tesla时代分离的纹理读取与ROP写入通道。
软件上，提出了统一内存地址空间，将Local、Shared、Global地址空间合并，开发者可通过统一指针进行访问，简化了编程模型，该特性在PTX 2.0中得以应用。

2.2 Fermi架构对应的产品

2.3 增加的CUDA操作

Fermi对应CUDA计算能力2.x，新增操作包括：

增强的线程同步函数：
```
int __syncthreads_count(int predicate);
int __syncthreads_and(int predicate);
int __syncthreads_or(int predicate);
```
这些函数在同步block内所有线程的同时，对每个线程的谓词值进行统计或逻辑运算。
系统级内存栅栏：__threadfence_system()，确保多GPU设备间的内存操作顺序全局可见。
增强的原子操作：支持共享内存上的64位整型原子操作，以及全局/共享内存上的32位浮点加法原子操作。
表面内存操作函数：新增一系列surf前缀函数，用于访问1D、2D表面内存。

3 Kepler 架构分析

Kepler架构于2012年4月发布，采用28nm工艺，芯片代码“GKXXX”。主要亮点：

动态并行：GPU线程可动态创建新的kernel。
Hyper-Q：提供最多32个硬件工作队列，提升多流并发度。
Warp Shuffle：warp内线程可直接通过寄存器交换数据，无需经过共享内存。
其他：增加Grid管理单元、GPU Boost、NVDEC/NVENC编解码器、无绑定纹理、四warp调度器、GPUDirect RDMA等。

3.1 架构特点

Kepler计算卡架构主体如下，包含15个SMX、6个内存控制器等。

图形卡架构（如GTX 680）则组织为多个GPC。

与Fermi的主要参数对比如下：

3.1.1 SMX特点
SMX的运算单元数量达到峰值（192个单精度CUDA Core），并首次引入独立的双精度单元（64个DP Unit）。

主要变化：

时钟调整：采用与GPU同频的1x时钟，相比前代的2x时钟降低了功耗。
四warp调度：SMX配备4个warp调度器，每个周期可发射4个warp的8条指令，且支持双精度与其它指令混合发射。
增加GMU：为实现动态并行，新增Grid管理单元(GMU)，用于创建和管理GPU内部发起的网格。

3.1.2 动态并行
允许GPU kernel在运行时创建并发射新的子kernel，无需CPU介入。优势在于可根据中间计算结果动态调整后续计算粒度，提升GPU利用率和任务独立性。

3.1.3 Hyper-Q
解决了Fermi单硬件工作队列的瓶颈。Kepler提供最多32个独立硬件队列，使来自不同CPU流或MPI任务的kernel能真正并发执行，提高了GPU利用率。

3.1.4 Warp Shuffle指令
提供__shfl()系列指令，允许warp内的32个线程直接通过寄存器交换数据，避免了通过共享内存中转的开销，极大提升了规约等操作的效率。

3.1.5 内存结构改进
新增独立的只读数据缓存，并为其开辟了专用通道。同时，纹理单元可被当作对象灵活使用，CUDA代码中可通过const __restrict修饰符让编译器自动使用该缓存。

3.1.6 GPUDirect RDMA
支持GPU之间通过PCIe或InfiniBand等网络直接访问对方显存，无需CPU内存中转，提升了多卡间数据交换效率。

3.2 Kepler架构对应的产品

3.3 新支持的CUDA操作

对应计算能力3.x，主要新特性：

使用只读数据缓存：

__global__ void saxpy(float x, float y, const float * __restrict input, float * output){
size_t offset = threadIdx.x + (blockIdx.x * blockDim.x);
// 编译器自动使用只读缓存加载input
output[offset] = (input[offset] * x) + y;
}

Warp Shuffle函数：实现warp内数据广播、规约等。

动态并行：示例子kernel创建。

__global__ void parent_kernel(int *data) {
if (threadIdx.x == 0) {
    child_kernel<<< 1, 256 >>>(data); // GPU内创建子kernel
    cudaDeviceSynchronize();
}
}