第2章：现代GPU架构概览：NVIDIA Ampere/Hopper与AMD RDNA/CDNA架构对比

各位同学，今天我们聊聊现代GPU的硬件架构。说实话，这个话题我每次讲都觉得特别有意思——你想想看，一块显卡里塞了几百亿个晶体管，它们怎么协同工作？为什么有的卡跑AI快，有的卡跑图形快？

我个人习惯把GPU架构比作一个大型工厂。NVIDIA和AMD就是两个不同的厂长，管理风格完全不同。咱们今天就来拆开看看，里面到底有什么门道。

2.1 NVIDIA Ampere与Hopper：从GA100到GH100

先看NVIDIA这边。Ampere架构（GA100）是2020年推出的，Hopper（GH100）是2022年的。我当年调试Ampere的驱动时，印象最深的就是它的SM（Streaming Multiprocessor）结构变化。

Ampere的SM内部结构：

• 每个SM包含4个处理块（Processing Block），每个块有16个INT32单元和16个FP32单元
• 关键变化：FP32和INT32可以同时执行了！之前Turing架构是共享的
• 每个SM有64KB的寄存器文件，分成4份，每份16KB
• 共享内存/L1缓存：最大164KB，可配置

Hopper的升级点：

• 引入了DPX指令——专门处理动态规划的。我当时看到这个指令集扩展，第一反应是「NVIDIA连这都想好了？」
• SM数量从Ampere的108个增加到132个（GH100）
• Transformer Engine：专门优化大模型训练，FP8精度支持
• 新的线程块集群（Thread Block Cluster）概念，跨SM协作更高效

2.2 AMD RDNA与CDNA：分道扬镳的两条路

AMD这边有意思了。他们干脆把图形和计算分成了两条产品线：RDNA管图形，CDNA管计算。

RDNA 3架构（RX 7000系列）：

• 每个CU（Compute Unit）包含2个SIMD单元，每个SIMD有32个ALU
• 引入了Infinity Cache——说白了就是一块超大L3缓存，最高96MB
• Wave32模式：和NVIDIA的Warp32类似，但AMD之前一直是Wave64
• 我测试过RDNA3的Shader执行效率，比RDNA2提升了约15%

CDNA 2架构（MI250/MI250X）：

• 每个CU有4个SIMD单元，比RDNA多一倍
• Matrix Core：对标NVIDIA的Tensor Core，专门做矩阵乘法
• 支持FP64双精度——RDNA砍掉了这个，CDNA保留
• Infinity Fabric互联：多卡通信延迟比PCIe低很多

2.3 SM vs CU：内部结构的深度对比

咱们把NVIDIA的SM和AMD的CU放在一起看看。说白了，它们都是「最小的独立计算单元」，但设计哲学完全不同。

特性	NVIDIA SM (Hopper)	AMD CU (RDNA3)
ALU数量	128 FP32 + 64 INT32	64 ALU (可执行FP32/INT32)
Warp/Wave大小	32线程 (Warp)	32线程 (Wave32) 或 64线程 (Wave64)
寄存器文件	64KB (每个SM)	128KB (每个CU)
共享内存	最大164KB (与L1共享)	64KB (专用)
特殊单元	Tensor Core, DPX	Matrix Core (CDNA)

嗯，这里要注意：NVIDIA的寄存器是「按SM分配」，AMD是「按CU分配」。这意味着什么？你想想看，NVIDIA的线程切换成本更低——因为寄存器就在SM内部。AMD的CU寄存器更大，但切换开销也更大。

2.4 内存层次结构：从全局到寄存器

内存层次结构，说白了就是「数据怎么从硬盘跑到计算单元里」。我当年优化一个矩阵乘法的kernel，就是因为没搞懂内存层次，性能差了10倍。

全局内存（Global Memory）：

• 就是显存，HBM2e或GDDR6X
• 带宽很大（Hopper H100有3.35TB/s），但延迟很高（几百个周期）
• 所有线程都能访问，但速度最慢

共享内存（Shared Memory）：

• 在SM/CU内部，所有线程共享
• 延迟低（几十个周期），容量小（几十KB）
• 我习惯用它做数据复用——比如矩阵分块计算

局部内存（Local Memory）：

• 每个线程私有的「伪寄存器」
• 实际上存在全局内存里，但编译器帮你管理
• 当寄存器不够用时，变量会被「溢出」到这里
• 性能杀手！我曾经遇到过寄存器溢出导致性能下降50%的情况

寄存器（Register）：

• 最快！一个周期就能访问
• 每个线程可用数量有限（NVIDIA一般是255个/线程）
• 编译器会尽量把变量分配到这里

2.5 实战对比：一个简单的向量加法

咱们用代码看看不同架构下的差异。假设我们要做向量加法：C = A + B。

// NVIDIA CUDA (Ampere/Hopper)
__global__ void vec_add(float* A, float* B, float* C, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        C[idx] = A[idx] + B[idx];
    }
}

// AMD HIP (RDNA/CDNA)
__global__ void vec_add(float* A, float* B, float* C, int N) {
    int idx = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
    if (idx < N) {
        C[idx] = A[idx] + B[idx];
    }
}

你看，代码几乎一模一样。但底层执行完全不同：

• NVIDIA上，32个线程组成一个Warp，一起执行
• AMD上，如果是Wave32模式，也是32个线程一起；如果是Wave64，就是64个
• 内存访问模式也不同——NVIDIA的L1和共享内存是统一的，AMD是分开的

我个人建议：写代码时别太在意架构差异，先把功能跑通。性能优化时再针对具体架构调整。毕竟，90%的性能问题都出在内存访问模式上，而不是计算单元。

2.6 小结与展望

今天咱们聊了NVIDIA和AMD的架构差异。说白了，NVIDIA更注重「通用性」——一个架构通吃图形和计算。AMD则选择「分而治之」——RDNA管图形，CDNA管计算。

下一章，我们会深入SM/CU内部的执行单元，看看指令是怎么被发射和执行的。到时候我会分享一个我调试Warp调度器的真实案例，保证让你大开眼界。