硬件架构基础：GPU/CPU计算单元、内存层次结构、PCIe带宽瓶颈、SM/Warp调度原理
各位同学，咱们今天聊点硬核的。做算子调优，说白了就是跟硬件打交道。你不摸清硬件的脾气，写出来的代码就是瞎蒙。我刚开始搞这行的时候，也踩过不少坑，今天把这些经验掰开了揉碎了讲给你听。
1. CPU vs GPU：两种不同的计算哲学
先说说CPU和GPU的区别。CPU像是个全能冠军，什么活都能干，但一次只能干几件事。GPU呢？它是个偏科生，只会做并行计算，但一做就是几千几万件事同时来。
我习惯这么理解：CPU是法拉利，跑得快但拉不了多少货；GPU是卡车队，跑得慢点但一次能拉一座山。你想想看，做矩阵乘法这种活，明显是卡车队更合适。
核心差异点：
- CPU：延迟优化，擅长复杂控制流，少量线程（几十个）
- GPU：吞吐优化，擅长数据并行，海量线程（几万个）
我在项目中遇到过一个问题：有人把CPU上的优化思路直接搬到GPU上，结果性能反而更差了。为什么？因为GPU的强项是隐藏延迟，不是降低延迟。这个观念不转过来，调优就无从谈起。
2. 内存层次结构：数据搬家的艺术
内存这玩意儿，离计算单元越近越快，但容量越小。说白了就是个金字塔结构。
我的经验：写算子时，脑子里要时刻想着数据在哪。L1缓存里？L2缓存里？还是HBM里？访问一次HBM的时间，够L1缓存跑几百次了。
2.1 CPU内存层次
| 层级 | 典型大小 | 延迟（周期） | 带宽 |
|---|---|---|---|
| L1缓存 | 32KB | 3-5 | 极高 |
| L2缓存 | 256KB-1MB | 10-20 | 高 |
| L3缓存 | 8-32MB | 30-50 | 中等 |
| DDR内存 | 16-512GB | 100-300 | 低 |
2.2 GPU内存层次
| 层级 | 典型大小 | 延迟（周期） | 带宽 |
|---|---|---|---|
| 寄存器 | 256KB/SM | 1 | 极高 |
| L1/共享内存 | 48-128KB/SM | 20-30 | 极高 |
| L2缓存 | 4-6MB | 100-200 | 高 |
| HBM | 16-80GB | 400-800 | 极高（1-2TB/s） |
嗯，这里要注意：GPU的HBM带宽看着吓人，但那是理论峰值。实际能跑到多少，取决于你的访存模式。我曾经调过一个算子，就因为数据没对齐，带宽利用率从80%掉到了20%。
3. PCIe带宽瓶颈：躲不开的痛
CPU和GPU之间靠PCIe总线通信。这玩意儿，说白了就是一根水管。水管再粗，也有上限。
避坑指南：我曾经以为PCIe 4.0 x16的32GB/s带宽够用了，结果发现频繁的kernel launch和内存拷贝，实际有效带宽连10GB/s都不到。为什么？因为每次传输都有开销，小数据包尤其吃亏。
PCIe带宽的几个关键点：
- 理论带宽 vs 有效带宽：通常只有60-70%
- 传输粒度：越大越好，尽量合并传输
- 双向传输：可以同时读写，但实际场景很少用满
我建议你在写算子时，尽量把数据留在GPU上。能一次传完的，别分十次。能复用数据的，别反复传。这个道理简单，但很多人就是做不到。
4. SM/Warp调度原理：GPU的微观世界
GPU的核心是SM（Streaming Multiprocessor）。每个SM里有一堆CUDA核心，还有调度器、寄存器、共享内存这些资源。
Warp是啥？就是32个线程组成的一个小组。这32个线程执行同一条指令，但处理不同的数据。这就是SIMT（单指令多线程）的精髓。
关键概念：
- 一个SM可以同时调度多个Warp
- Warp切换是零开销的（硬件自动完成）
- Warp内的线程必须执行同一条指令（分支发散会串行化）
我刚开始学的时候，总觉得Warp调度很神秘。后来发现，说白了就是：GPU用大量线程来隐藏延迟。一个Warp在等数据，就切到另一个Warp执行。只要你有足够的Warp，计算单元就永远不会闲着。
4.1 分支发散：性能杀手
Warp内的32个线程，如果走不同的分支，会发生什么？
// 坏例子：分支发散
if (threadIdx.x % 2 == 0) {
// 偶数线程执行这里
a[threadIdx.x] = b[threadIdx.x] * 2;
} else {
// 奇数线程执行这里
a[threadIdx.x] = b[threadIdx.x] + 1;
}
// 实际执行：先执行偶数分支（奇数线程被屏蔽），再执行奇数分支（偶数线程被屏蔽）
// 效率直接砍半！我曾经在一个图像处理算子中，因为没注意分支发散，性能只有预期的40%。后来改成用位运算和条件赋值，才把性能拉回来。
4.2 共享内存与Bank Conflict
共享内存是GPU上最快的可编程存储。但它有个坑：Bank Conflict。
共享内存被分成32个Bank（对应32个线程）。如果同一个Warp里的多个线程访问同一个Bank的不同地址，就会发生冲突，访问会串行化。
我的习惯：写共享内存代码时，先画个访问模式图。看看每个线程访问的是哪个Bank。如果发现冲突，就加padding或者调整数据布局。
5. 知识体系总览
下面这张图，是我自己总结的算子调优知识体系。你把它记在脑子里，调优时就知道该往哪个方向使劲了。
这张图把今天讲的内容串起来了。你从计算单元出发，理解CPU和GPU的差异；然后看内存层次，知道数据在哪、怎么搬；再考虑PCIe这个瓶颈；最后深入到SM内部，理解Warp怎么调度、怎么避免分支发散和Bank Conflict。
嗯，今天就先聊到这儿。这些基础概念，你回去好好消化一下。下一节我们开始讲具体的算子实现，到时候这些知识全都会用上。