GPU 硬件架构深度解析
理解 GPU 硬件是写出高性能 CUDA 代码的前提。本文从芯片级别拆解 NVIDIA GPU 的物理结构，揭示并行计算能力的硬件来源�?
GPU 整体架构层次
GPU Die
�?
├── GPC (Graphics Processing Cluster) × N
�? ├── TPC (Texture Processing Cluster) × M
�? �? └── SM (Streaming Multiprocessor) × 2
�? �? ├── CUDA Core (FP32) × 128
�? �? ├── Tensor Core × 4
�? �? ├── RT Core × 1 (Turing+)
�? �? ├── LD/ST Unit × 32
�? �? ├── SFU (Special Function Unit) × 16
�? �? ├── Warp Scheduler × 4
�? �? ├── Register File (256KB)
�? �? ├── L1 Cache / Shared Memory (128KB)
�? �? └── Texture Unit
�? �?
�? └── Raster Engine
�?
├── L2 Cache (共享，全芯片)
├── Memory Controller × N
└── HBM / GDDR 显存SM �?Streaming Multiprocessor
SM �?GPU �?基本计算单元*，所�?CUDA Kernel 都在 SM 上执行�?
SM 内部结构（以 Ampere A100 为例�?
| 组件 | 数量 | 说明 |
|---|---|---|
| FP32 CUDA Core | 128 | 单精度浮�?整数运算 |
| FP64 Core | 64 | 双精度浮点运�? |
| Tensor Core (3rd Gen) | 4 | 矩阵乘法加速单�? |
| LD/ST Unit | 32 | 内存加载/存储单元 |
| SFU | 16 | 超越函数（sin/cos/sqrt/rcp�? |
| Warp Scheduler | 4 | 每周期各发射1条指�? |
| Register File | 256 KB | �?SM 最�?65536 �?32-bit 寄存�? |
| L1 Cache / Shared Mem | 192 KB | 可配置分配比�? |
A100 整体规格
A100 SXM4
├── SM 数量�?08
├── CUDA Core�?08 × 128 = 13,824
├── Tensor Core�?08 × 4 = 432
├── FP32 峰值：19.5 TFLOPS
├── TF32 Tensor Core�?56 TFLOPS
├── BF16 Tensor Core�?12 TFLOPS
├── INT8 Tensor Core�?24 TOPS
├── 显存�?0GB HBM2e
└── 显存带宽�? TB/sWarp �?执行的基本单�?
什么是 Warp
**Warp �?GPU 调度和执行的最小单�?*，由 **32 个线�?*组成�?
Block (例如 256 个线�?
├── Warp 0：Thread 0-31
├── Warp 1：Thread 32-63
├── Warp 2：Thread 64-95
├── ...
└── Warp 7：Thread 224-255SIMT 执行模型
Warp �?32 个线�?同时执行同一条指�?，但各自操作不同数据�?
cpp
// 每个线程执行相同指令，但 threadIdx.x 不同
__global__ void add(float* a, float* b, float* c) {
int i = threadIdx.x + blockIdx.x * blockDim.x;
c[i] = a[i] + b[i]; // 32个线程同时执行这�?
}Warp Divergence（分支分歧）
性能陷阱
�?Warp 内线程走不同分支时，GPU 必须串行执行每个分支，效率大幅下降�?
cpp
// 坏：Warp 内线程分�?
if (threadIdx.x % 2 == 0) {
// 偶数线程执行
result = a * b;
} else {
// 奇数线程执行（与偶数线程串行！）
result = a + b;
}
// 好：�?Warp 对齐的分支（32的倍数�?
if (threadIdx.x / 32 == 0) {
// Warp 0 的所有线程走这里
result = a * b;
} else {
// 其他 Warp 走这�?
result = a + b;
}Warp 占用率（Occupancy�?
*Occupancy = 活跃 Warp �?/ SM 最�?Warp �?
SM 通过在多�?Warp 间切换来隐藏内存延迟�?
时间轴：
Warp A: [计算] [等待内存] ........... [计算]
Warp B: [计算] [等待内存] ... [计算]
Warp C: [计算] [等待内存] [计算]
�?�?A 等待时，调度 B/C 执行影响 Occupancy 的因素：
| 资源 | 限制 |
|---|---|
| 寄存器数�? | �?SM 65536 个，寄存器越多，活跃 Block 越少 |
| 共享内存 | �?SM 固定大小，用得越多，活跃 Block 越少 |
| Block 大小 | 太小导致 Warp 不足，太大限制并�?Block �? |
Tensor Core
Tensor Core 是专�?*矩阵乘法**设计的硬件加速单元，是深度学习性能的核心来源�?
工作原理
每个 Tensor Core 每个时钟周期执行一�?**4×4 矩阵乘加（MMA�?*�?
D = A × B + C
其中 A(4×4), B(4×4), C(4×4), D(4×4)各代 Tensor Core 对比
| 架构 | 代数 | 支持精度 | 峰值（单卡�? |
|---|---|---|---|
| Volta (V100) | 1st | FP16 | 125 TFLOPS |
| Turing (T4) | 2nd | FP16, INT8, INT4 | 65 TOPS |
| Ampere (A100) | 3rd | FP16, BF16, TF32, INT8 | 312 TFLOPS |
| Hopper (H100) | 4th | FP8, FP16, BF16, TF32 | 989 TFLOPS |
| Blackwell (B200) | 5th | FP4, FP8, FP16, BF16 | 4.5 PFLOPS |
TF32 精度格式
Ampere 引入�?TF32（TensorFloat-32）：
FP32: 1位符�?+ 8位指�?+ 23位尾�?= 32�?
TF32: 1位符�?+ 8位指�?+ 10位尾�?= 19�?
FP16: 1位符�?+ 5位指�?+ 10位尾�?= 16�?
TF32 = FP32 的指数范�?+ FP16 的精�?
�?对深度学习训练几乎无精度损失，速度提升 10x内存系统架构
访问延迟（从快到慢）�?
寄存�? ~1 cycle 每线程独有，最�?
L1 Cache ~20 cycles �?SM 共享�?28-192KB�?
共享内存 ~20 cycles �?SM 共享，程序员控制
L2 Cache ~200 cycles 全芯片共享（A100: 40MB�?
全局内存(HBM) ~600 cycles 所�?SM 共享，容量最�?内存带宽对比
| 内存类型 | A100 带宽 | H100 带宽 |
|---|---|---|
| HBM2e/HBM3 | 2 TB/s | 3.35 TB/s |
| L2 Cache | ~12 TB/s | ~60 TB/s |
| 共享内存 | ~19 TB/s/SM | ~33 TB/s/SM |
NVLink �?GPU 互联
当多�?GPU 协同工作时，NVLink 提供高带宽的 GPU 间直连：
| 版本 | 单链路带�? | 链路�? | 总带�? |
|---|---|---|---|
| NVLink 2.0 (V100) | 25 GB/s × 2 | 6 | 300 GB/s |
| NVLink 3.0 (A100) | 25 GB/s × 2 | 12 | 600 GB/s |
| NVLink 4.0 (H100) | 25 GB/s × 2 | 18 | 900 GB/s |
对比 PCIe 5.0 x16 �?~64 GB/s，NVLink 带宽高出一个数量级�?
关键性能指标
Roofline 模型
性能 = min(峰值算�? 内存带宽 × 算术强度)
算术强度 = FLOP / Byte（每字节内存访问对应的浮点运算数�?
峰值算�?
�?
性能 �? ╱‾‾‾‾‾‾‾‾‾‾‾‾‾‾�?
(FLOPS) �? �? 内存带宽受限 �?计算受限
�? �?
�? �?
└──────────────────────────
算术强度 (FLOP/Byte)优化方向
- 算术强度�?�?优化内存访问（合并访问、使用共享内存）
- 算术强度高但未达峰�?�?提高 Occupancy、减少分支分�?