第三章：GPU硬件架构详解——SM、CUDA Core、Warp调度与内存层次
各位同学，今天我们来聊聊GPU的“骨架”。
说实话，我刚开始接触CUDA时，最头疼的就是搞不清硬件到底长什么样。你写个kernel，它怎么就跑起来了？为什么有些代码快得飞起，有些却慢得像蜗牛？
嗯，答案全在硬件架构里。这一章，我们就把GPU的“五脏六腑”拆开看看。
3.1 流多处理器（SM）——GPU的“计算单元”
先问个问题：GPU和CPU最大的区别在哪？
CPU是“全能选手”，什么活都能干，但一次只能干几件。GPU呢？它是“人海战术”的典范——成千上万个简单计算单元同时干活。
这个“简单计算单元”的集合，就是SM（Streaming Multiprocessor，流多处理器）。
核心概念：SM是GPU内部最基本的独立计算单元。一个GPU芯片里，通常有几十到上百个SM。每个SM都有自己的控制单元、执行单元、缓存和寄存器文件。
我习惯把SM比作一个“小工厂”。每个工厂里有很多工人（CUDA Core），有仓库（共享内存），有调度员（Warp Scheduler）。你写的CUDA程序，最终就是被分配到这些工厂里执行的。
举个例子，NVIDIA的A100 GPU有108个SM。每个SM内部有64个FP32 CUDA Core。算一下：108 × 64 = 6912个核心。这就是为什么GPU能同时处理那么多线程。
个人经验：我在做大规模并行计算时，经常需要估算SM数量。比如你要处理100万个数据点，如果SM太少，线程块分配就会很拥挤。我一般会先查一下目标GPU的SM数量，再决定线程块的大小。
3.2 CUDA Core——真正的“打工人”
SM里的CUDA Core，才是真正干活的。
每个CUDA Core本质上是一个浮点运算单元，能执行整数和浮点数的加减乘除。但注意，它不像CPU核心那样有复杂的乱序执行、分支预测功能。说白了，它就是个“傻大个”，只会按指令干活。
一个典型的CUDA Core包含：
- 整数运算单元（INT ALU）：处理整数加减、位运算等
- 单精度浮点单元（FP32）：处理32位浮点数运算
- 双精度浮点单元（FP64）：处理64位浮点数（数量通常只有FP32的1/2或1/4）
- 张量核心（Tensor Core）：专门为AI矩阵运算设计的加速单元
| 核心类型 | 用途 | 性能特点 |
|---|---|---|
| INT32 | 整数运算、地址计算 | 快，但精度有限 |
| FP32 | 通用浮点计算 | 最常用，平衡精度和速度 |
| FP64 | 科学计算、高精度需求 | 慢，数量少 |
| Tensor Core | 矩阵乘法、深度学习 | 极快，但需要特定数据格式 |
避坑指南：我曾经有个项目，用FP64做物理模拟，结果性能惨不忍睹。后来才发现，我的GPU（RTX 3080）FP64核心只有FP32的1/64。换成FP32后，速度提升了30倍。所以，不是所有核心都生而平等。
3.3 Warp调度——GPU的“指挥艺术”
现在问题来了：SM里有那么多CUDA Core，怎么让它们高效工作？
答案是Warp。
Warp是GPU调度的基本单位。一个Warp包含32个线程。这32个线程在同一个SM上执行相同的指令，但处理不同的数据。这就是所谓的SIMT（单指令多线程）模型。
每个SM里有多个Warp Scheduler（调度器）。它的工作就是：
- 从就绪的Warp中选一个
- 发射一条指令给CUDA Core
- 切换到下一个Warp
- 重复以上步骤
为什么要这样？因为内存访问很慢。当一个Warp在等数据时，调度器可以立刻切换到另一个Warp，让CUDA Core不闲着。这叫延迟隐藏。
关键点：Warp内的32个线程必须执行同一条指令。如果遇到分支（if-else），部分线程会“休眠”，等另一部分执行完再唤醒。这就是Warp Divergence（分支发散），会严重降低性能。
我记得有一次优化一个图像处理算法，发现性能始终上不去。后来用nvprof一查，Warp Divergence高达40%。原来是我在循环里用了if (threadIdx.x % 2 == 0)这样的分支。改成按Warp对齐后，性能直接翻倍。
建议：写CUDA代码时，尽量让Warp内的32个线程走相同的路径。如果必须分支，让分支条件基于warpId（线程ID除以32），而不是基于单个线程ID。
3.4 内存层次结构——GPU的“存储迷宫”
GPU的内存系统，是性能优化的重中之重。我见过太多人，代码写得漂亮，但内存访问一塌糊涂，性能惨不忍睹。
GPU的内存从快到慢、从近到远，分为这几个层次：
3.4.1 寄存器（Register）
最快，每个线程私有。容量极小，每个线程最多255个32位寄存器。如果你用了太多局部变量，寄存器不够用，数据会被“溢出”到本地内存（其实在全局内存里），速度骤降。
3.4.2 共享内存（Shared Memory）
在SM内部，所有线程块内的线程可以共享。速度很快，但容量有限（通常几十KB到上百KB）。
共享内存是手动管理的缓存。你可以在__shared__关键字声明它。我经常用它来做数据复用，比如矩阵分块计算。
// 共享内存示例：矩阵分块
__global__ void matmul_shared(float *A, float *B, float *C, int N) {
__shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
__shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
int bx = blockIdx.x, by = blockIdx.y;
int tx = threadIdx.x, ty = threadIdx.y;
float sum = 0.0f;
for (int k = 0; k < N / BLOCK_SIZE; ++k) {
// 加载数据到共享内存
As[ty][tx] = A[by * BLOCK_SIZE + ty][k * BLOCK_SIZE + tx];
Bs[ty][tx] = B[k * BLOCK_SIZE + ty][bx * BLOCK_SIZE + tx];
__syncthreads(); // 同步，确保所有线程加载完成
// 计算
for (int i = 0; i < BLOCK_SIZE; ++i)
sum += As[ty][i] * Bs[i][tx];
__syncthreads();
}
C[by * BLOCK_SIZE + ty][bx * BLOCK_SIZE + tx] = sum;
}3.4.3 L1/L2缓存
L1缓存和共享内存共用同一块物理存储（通常可配置）。L2缓存是所有SM共享的，容量更大（几MB到几十MB）。
缓存是自动管理的，但你可以通过__ldg()等指令提示编译器使用只读缓存。
3.4.4 全局内存（Global Memory）
最大的存储空间（几GB到几十GB），但最慢。延迟通常400-800个时钟周期。
访问全局内存时，要特别注意合并访问（Coalesced Access）。说白了，就是让同一个Warp的32个线程访问连续的内存地址。这样GPU可以把多次访问合并成一次大的内存事务。
合并访问示例：
好的访问模式：data[threadIdx.x]（连续）
坏的访问模式：data[threadIdx.x * 32]（跨步）
坏的访问模式：data[blockIdx.x * blockDim.x + threadIdx.x]（如果blockDim.x不是32的倍数）
3.4.5 常量内存和纹理内存
常量内存：只读，有缓存，适合存储系数、查找表等。所有线程访问同一地址时很快。
纹理内存：专门为图像处理设计，支持二维/三维空间局部性访问。
| 内存类型 | 位置 | 访问速度 | 作用域 | 生命周期 |
|---|---|---|---|---|
| 寄存器 | SM内部 | 1个周期 | 单个线程 | 线程内 |
| 共享内存 | SM内部 | ~5个周期 | 线程块内 | 线程块内 |
| L1缓存 | SM内部 | ~10个周期 | SM内 | 自动管理 |
| L2缓存 | 芯片共享 | ~100个周期 | 全局 | 自动管理 |
| 全局内存 | 显存 | 400-800个周期 | 全局 | 程序运行期间 |
| 常量内存 | 显存（有缓存） | ~100个周期（缓存命中） | 全局 | 程序运行期间 |
我曾经踩过的坑：有次做粒子模拟，每个粒子需要访问一个全局数组。我图省事，直接让每个线程独立访问。结果因为访问模式不连续，内存带宽利用率只有5%。后来改成共享内存做缓存，先让线程块内的线程协作加载连续数据，再各自计算。性能提升了15倍。
记住：全局内存访问，能合并就合并，能缓存就缓存。
3.5 总结与思考
好了，这一章的内容就到这里。我们来捋一捋：
- SM是GPU的基本计算单元，里面有CUDA Core、调度器、缓存等
- CUDA Core是执行算术运算的“工人”，不同类型核心性能差异巨大
- Warp是调度单位，32个线程一起执行，要避免分支发散
- 内存层次从快到慢：寄存器 → 共享内存 → L1/L2缓存 → 全局内存
我个人觉得，理解这些硬件细节，是写出高性能CUDA代码的必经之路。你不需要记住每个GPU的精确参数，但一定要明白：你的代码最终是在这些硬件上运行的，它的性能受限于硬件的特性。
下一章，我们会深入CUDA编程模型，看看如何用代码控制这些硬件资源。到时候，我会分享一些我实际项目中用过的优化技巧。
嗯，今天就到这里。有问题随时交流。