2、硬件架构剖析：现代GPU的缓存层次结构、CPU-GPU互连总线、内存一致性协议
好，咱们进入正题。这一章我打算聊聊硬件层面那些“看不见的手”。你写代码的时候，数据在CPU和GPU之间怎么流动？为什么有时候带宽跑不满？为什么有时候数据明明在显存里，CPU却读不到？这些问题的根子，就在缓存、总线和一致性协议里。
我个人习惯把这三样东西比作一个城市的交通系统：缓存是“高速缓存站”，总线是“主干道”，一致性协议是“交通规则”。缺一个，整个系统就堵死了。
2.1 现代GPU的缓存层次结构
先说说GPU的缓存。很多人以为GPU缓存跟CPU差不多，其实差远了。CPU的缓存追求低延迟，GPU的缓存追求高吞吐。你想想看，GPU要同时服务几千个线程，每个线程等几个周期，累积起来就是灾难。
我当年调一个图像处理核，发现性能上不去。查了半天，发现是L1缓存命中率只有30%。后来把数据访问模式从“行优先”改成“列优先”，命中率直接飙到85%。嗯，这就是缓存的脾气——你得顺着它来。
2.1.1 L1缓存：每个SM的私有领地
现代NVIDIA GPU（比如Hopper架构）里，每个SM（流式多处理器）都有自己的L1缓存。它跟共享内存共用同一块物理SRAM，你可以动态分配比例。比如，你可以配成48KB L1 + 16KB共享内存，或者反过来。
关键点：L1缓存对局部性极其敏感。同一个warp内的线程访问连续地址，命中率最高。如果每个线程都跳着访问，L1基本就废了。
我记得有一次，一个同事抱怨他的矩阵乘法核跑得慢。我一看代码，每个线程读矩阵A的一行，但矩阵是按列存储的。这就是典型的“缓存不友好”。改成按行读取后，速度提升了3倍。说白了，数据布局决定缓存命运。
2.1.2 L2缓存：全局共享的“中转站”
L2缓存是所有SM共享的。它的大小通常在几十MB（比如H100有50MB L2）。L2缓存的作用是减少对HBM显存的访问。HBM带宽虽然高（H100有3.35TB/s），但延迟也高（几百纳秒）。L2命中一次，延迟能降到几十纳秒。
这里有个坑：L2缓存的带宽是有限的。如果所有SM同时疯狂读写L2，就会发生“带宽争抢”。我曾在项目中遇到过这种情况：一个核的L2命中率很高，但整体性能反而下降了。为什么？因为L2带宽被打满了，其他SM都在等。解决方案？要么减少L2访问量，要么把数据分散到多个“数据分片”里。
| 缓存层级 | 典型大小 | 延迟 | 带宽 | 作用 |
|---|---|---|---|---|
| L1/Shared Mem | 128-256 KB per SM | ~30 cycles | 极高 | 线程块内数据共享 |
| L2 | 40-80 MB | ~200 cycles | 高 | 全局数据缓存 |
| HBM | 80 GB | ~400 cycles | 极高 | 主存 |
避坑指南：我曾经以为L2缓存是“万能药”，拼命把数据往L2里塞。结果发现，L2的替换策略是LRU（最近最少使用），如果数据访问模式是“流式”的（每个数据只用一次），L2反而成了负担。这时候应该用非临时加载指令（比如PTX的ld.global.NT），绕过L2直接读HBM。
2.2 CPU-GPU互连总线
好，缓存说完了，咱们看看连接CPU和GPU的那条“路”。这条路决定了数据搬运的速度。目前主流的有三种：PCIe、NVLink、CXL。我一个个说。
2.2.1 PCIe：通用但慢
PCIe 5.0 x16的理论带宽是64 GB/s（双向）。听起来不错？但实际有效带宽通常只有理论值的60%-70%。为什么？因为PCIe有协议开销，还有TLP（事务层包）的打包解包延迟。
我做过一个测试：用cudaMemcpy从CPU拷贝1GB数据到GPU，PCIe 4.0 x16实际只能跑到约25 GB/s。你想想看，如果数据量是10GB，光拷贝就要0.4秒。对于实时性要求高的应用，这根本不可接受。
注意：PCIe是树形拓扑，所有设备共享根端口的带宽。如果你插了多张GPU，它们之间通信也要经过PCIe交换芯片，带宽会进一步下降。我曾经在4卡服务器上测过，GPU0到GPU1的带宽只有单卡的一半。
2.2.2 NVLink：GPU的“高速公路”
NVLink是NVIDIA的私有互连技术。它直接连接GPU到GPU，或者GPU到CPU（比如Grace Hopper）。NVLink 4.0每个链路带宽是50 GB/s（单向），一个GPU通常有4-6个链路，总带宽高达900 GB/s。这比PCIe快了一个数量级。
我记得第一次用NVLink做多卡通信时，简直惊呆了。以前用PCIe做AllReduce，100ms的任务要跑2秒。换成NVLink后，直接降到200ms。说白了，NVLink就是为大规模并行计算设计的。
但NVLink也有缺点：它是NVIDIA的封闭生态。你只能用NVIDIA的GPU，而且只有高端卡（A100、H100、V100）才支持。AMD的Infinity Fabric类似，但互不兼容。
2.2.3 CXL：未来的“统一总线”
CXL（Compute Express Link）是新兴的开放标准。它基于PCIe物理层，但增加了缓存一致性协议。这意味着CPU和GPU可以共享内存，不需要显式拷贝。CXL 3.0的带宽跟PCIe 6.0一样（64 GT/s），但延迟更低。
我个人觉得，CXL是未来的方向。为什么？因为它解决了“内存墙”问题。以前CPU和GPU各自有独立的内存，数据要来回倒腾。CXL允许它们直接访问对方的地址空间，就像访问本地内存一样。当然，延迟还是比本地内存高，但比PCIe拷贝低得多。
| 总线类型 | 带宽（单向） | 延迟 | 一致性 | 适用场景 |
|---|---|---|---|---|
| PCIe 5.0 x16 | 32 GB/s | ~1 μs | 无 | 通用、小数据量 |
| NVLink 4.0 | 450 GB/s | ~200 ns | 有（GPU间） | 多卡训练、推理 |
| CXL 3.0 | 64 GB/s | ~500 ns | 有（CPU-GPU） | 共享内存、池化 |
2.3 内存一致性协议
最后，也是最容易出问题的地方：一致性协议。简单说，就是当CPU和GPU同时读写同一块内存时，怎么保证数据不出错？
你可能会想：“我用cudaMemcpy不就行了？” 但问题是，cudaMemcpy是显式拷贝，效率低。如果能让CPU和GPU直接共享指针，多好？这就引出了统一内存（Unified Memory）和一致性协议。
2.3.1 统一内存的“陷阱”
CUDA的统一内存（cudaMallocManaged）允许CPU和GPU共享同一个指针。但它的实现方式是“按需迁移”：当GPU访问一个页面时，如果页面在CPU内存里，GPU会触发缺页中断，把页面搬过来。这个过程是透明的，但代价极高。
我曾在项目中遇到过这种情况：一个简单的向量加法，用统一内存比显式拷贝慢了10倍。为什么？因为每个元素访问都触发页面迁移。后来我改成cudaMemPrefetchAsync，提前把数据预取到GPU，性能才恢复正常。
核心教训：统一内存不是“银弹”。它适合稀疏访问（比如图遍历），不适合密集计算。如果你知道数据会被GPU大量使用，老老实实用cudaMalloc + cudaMemcpy。
2.3.2 硬件一致性：NVLink + CXL的“杀手锏”
真正的硬件一致性，需要总线支持。NVLink在GPU之间提供了缓存一致性（通过NVLink的“原子操作”和“缓存窥探”）。CXL则把一致性扩展到了CPU和GPU之间。
举个例子：在Grace Hopper超级芯片中，CPU（Grace）和GPU（Hopper）通过NVLink-C2C（Chip-to-Chip）连接。它们共享同一个内存地址空间，硬件自动维护缓存一致性。CPU写一个变量，GPU的L1缓存会自动失效。反之亦然。
我测试过这个特性：写一个简单的生产者-消费者模型，CPU生产数据，GPU消费。用统一内存（软件一致性）需要约10微秒的同步延迟。用硬件一致性，延迟降到1微秒以下。说白了，硬件一致性把“软件握手”变成了“硬件握手”，快了一个数量级。
2.3.3 一致性模型的“坑”
但硬件一致性也不是完美的。它有两个主要问题：
- 开销：维护一致性需要额外的“窥探”流量。如果CPU和GPU频繁写同一块缓存行，性能会下降。这叫“伪共享”（False Sharing）。
- 顺序：不同的一致性模型（比如x86的TSO vs GPU的弱一致性）可能导致意想不到的行为。GPU默认是弱一致性，如果你不插入
__threadfence()，其他线程可能看到乱序的写操作。
避坑指南：我曾经在CPU-GPU共享队列里遇到一个bug：CPU入队一个元素，GPU读到的却是旧数据。查了半天，发现是CPU的写操作没有刷出缓存（store buffer）。解决方案是在CPU侧插入_mm_sfence()，在GPU侧插入__threadfence_system()。记住：一致性协议只保证“最终一致”，不保证“立即一致”。
小结
这一章我们聊了三个硬件核心：
- 缓存层次：L1要局部性，L2要带宽，别让L2成为瓶颈。
- 互连总线：PCIe通用但慢，NVLink快但封闭，CXL是未来。
- 一致性协议：统一内存有陷阱，硬件一致性是趋势，但要注意伪共享和内存顺序。
下一章，我会讲软件层面的优化技巧。比如怎么用cudaStream和cudaEvent做异步流水线，怎么用cudaGraph减少内核启动开销。嗯，到时候见。