1. GPU架构基础：从图形卡到通用计算
大家好，我是你们这期课程的讲师。今天咱们聊聊GPU的硬件架构。说实话，我刚开始接触GPU时，也以为它就是个跑图形的卡。后来做高性能计算项目，才真正理解它的设计哲学。
GPU和CPU最大的区别是什么？CPU是「精打细算」的，它擅长处理复杂的逻辑分支。而GPU是「人多力量大」，它靠大量的计算单元堆出性能。你想想看，一个CPU可能只有8个核，而一个GPU有几千个CUDA Core。这就是根本差异。
1.1 从图形卡到通用计算
早期的GPU，说白了就是专门画图的。它负责把三角形、纹理、像素这些东西渲染到屏幕上。那时候的程序员想用它做别的计算？没门。
转折点出现在2006年。NVIDIA推出了CUDA架构。我记得当时看到这个技术，第一反应是：「这玩意儿能用来做科学计算？」后来我参与了一个分子动力学模拟项目，用GPU加速后，计算速度提升了40倍。嗯，从那以后我就彻底信了。
通用计算GPU（GPGPU）的核心思想很简单：把数据并行的问题扔给GPU。比如矩阵乘法、图像处理、深度学习训练——这些任务天然适合GPU。
关键点：GPU不是万能的。它擅长「单指令多数据」（SIMD）模式。如果你的程序分支多、数据依赖强，GPU反而跑不过CPU。我在项目中见过有人硬把串行逻辑塞进GPU，结果性能还不如CPU。
1.2 SM：GPU的核心计算单元
SM的全称是Streaming Multiprocessor，流式多处理器。它是GPU的基本计算单元。你可以把SM想象成一个「小型的计算集群」。每个SM内部有自己的寄存器、共享内存、调度器。
一个GPU有多少个SM？这取决于具体型号。比如NVIDIA的A100有108个SM，而RTX 3090有82个SM。SM越多，并行能力越强。
我个人习惯把SM比作「工厂车间」。每个车间有自己的工人（CUDA Core）、工具（寄存器）和仓库（共享内存）。车间之间可以协作，但主要靠各自干活。
避坑指南：我曾经犯过一个错误——以为SM越多性能就一定越好。后来发现，如果程序的内存访问模式很差，SM再多也白搭。因为数据喂不饱计算单元。这就像给100个厨师每人一把小刀，但食材供应跟不上，照样做不出菜。
1.3 CUDA Core：真正的执行者
CUDA Core是SM里的基本计算单元。每个CUDA Core可以执行一个浮点或整数运算。注意，这里的「Core」和CPU的Core不是一回事。CPU的Core是独立的，而CUDA Core是SM的一部分，它们共享控制逻辑和缓存。
举个例子：一个SM通常包含64个或128个CUDA Core。这些Core以32个为一组，称为一个Warp。Warp是GPU调度和执行的基本单位。为什么是32？这是NVIDIA的设计选择。我猜他们做过大量测试，发现32是最优的。
这里有个容易混淆的地方：CUDA Core的数量并不直接等于性能。比如RTX 3080有8704个CUDA Core，而A100有6912个。但A100的浮点性能反而更高。为什么？因为A100的每个CUDA Core支持更多的并发指令，而且内存带宽更大。
注意：别被CUDA Core的数量忽悠了。我见过一些厂商拿这个数字做营销。实际性能要看：时钟频率、内存带宽、缓存大小、以及你的程序能不能充分利用这些资源。
1.4 Memory Hierarchy：存储层次
GPU的存储层次是性能优化的关键。说白了，就是数据从哪来、到哪去、速度有多快。我把它分成几个层级：
| 层级 | 容量 | 带宽 | 延迟 | 作用 |
|---|---|---|---|---|
| 寄存器 | 每个线程几十个 | 极高 | 1个周期 | 存储局部变量 |
| 共享内存 | 每个SM几十KB | 高 | 几个周期 | 线程间通信 |
| L1缓存 | 每个SM几十KB | 高 | 十几个周期 | 缓存全局数据 |
| L2缓存 | 整个GPU几MB | 中等 | 几十个周期 | 跨SM共享 |
| 全局内存 | 几GB到几十GB | 低 | 几百个周期 | 主存储 |
看到这个表格，你可能会问：为什么全局内存这么慢？因为它是通过DRAM实现的，物理上离计算单元远。而寄存器就在SM内部，所以快。
我在优化一个图像处理算法时，发现程序大部分时间都花在全局内存访问上。后来我把频繁访问的数据挪到共享内存里，性能提升了3倍。嗯，这就是存储层次的价值。
1.5 实战：如何利用存储层次
光讲理论没意思。咱们看个实际例子。假设你要做矩阵乘法：C = A * B。如果直接写，每个线程从全局内存读数据，那性能会惨不忍睹。
优化思路是什么？把A和B的子块加载到共享内存里。这样每个线程只需要从共享内存读数据，速度就快多了。代码大概长这样：
// 每个线程块处理一个子块
__shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
__shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
// 从全局内存加载到共享内存
As[tx][ty] = A[row][ty];
Bs[tx][ty] = B[tx][col];
__syncthreads();
// 从共享内存计算
for (int k = 0; k < BLOCK_SIZE; ++k) {
Cvalue += As[tx][k] * Bs[k][ty];
}这段代码里，__syncthreads()是同步点。它保证所有线程都加载完数据后，才开始计算。我刚开始写CUDA时，经常忘记加这个同步，结果算出来的结果全是错的。后来养成了习惯：每次加载完共享内存，立刻加同步。
个人经验：共享内存的大小是有限的。比如每个SM只有48KB。如果你的子块太大，放不下怎么办？那就分多次加载。这叫「分块」（tiling）技术。我在一个深度学习项目中，用分块技术把卷积运算加速了5倍。
1.6 总结与思考
这一章我们聊了GPU的硬件架构。核心就三点：
- SM是计算单元，相当于工厂车间
- CUDA Core是执行单元，相当于工人
- 存储层次决定了数据访问速度
你可能会问：这些知识对性能优化有什么用？用处大了。比如你发现程序跑得慢，先看是不是全局内存访问太多。如果是，就考虑用共享内存或寄存器。再比如，你的Warp发散严重，那就调整线程布局。
下一章，我们会深入讨论内存带宽的瓶颈。我会用实际案例告诉你，为什么有时候GPU利用率很高，但性能就是上不去。嗯，敬请期待。