2. 缓存基础：CPU缓存与GPU缓存的异同，为什么GPU需要缓存
好，咱们进入第二讲。缓存这东西，说白了就是「快慢之间的缓冲带」。CPU有缓存，GPU也有缓存，但两者的脾气秉性完全不同。我刚开始接触GPU编程时，总拿CPU那套缓存思维去套GPU，结果性能惨不忍睹。后来才明白，这两兄弟虽然都叫缓存，但设计哲学天差地别。
2.1 CPU缓存：为「低延迟」而生
CPU缓存的设计目标很单纯——让单线程跑得更快。你想想看，CPU的主频动辄3-4GHz，而内存的访问延迟还在100纳秒级别。这中间的差距，全靠缓存来填。
典型的CPU缓存分三级：L1、L2、L3。L1最快也最小，通常32KB左右，延迟只有几个时钟周期。L3最大，几十MB，延迟几十个周期。我做过一个项目，把热点数据从L3挪到L1，性能直接翻倍。嗯，这就是缓存的魔力。
CPU缓存的核心特征：
- 低延迟优先：L1缓存延迟仅3-5个时钟周期
- 一致性协议强：MESI协议保证所有核看到的数据一致
- 容量小：L1通常32KB，L3也就几十MB
- 硬件预取：自动识别访问模式并提前加载数据
我曾经踩过一个坑：在多线程程序中，两个核频繁修改同一个缓存行，结果触发了「伪共享」问题。性能从100%直接掉到20%。后来用缓存行对齐才解决。这个教训让我深刻理解了CPU缓存一致性协议的代价。
2.2 GPU缓存：为「高吞吐」而生
GPU的缓存设计思路完全不同。GPU要处理的是成千上万个线程，它不在乎单个线程的延迟，只在乎整体吞吐量。说白了，CPU是法拉利，追求单次加速；GPU是公交车，追求一次拉很多人。
GPU的缓存层次更简单：L1缓存（也叫Shared Memory）和L2缓存。L1是每个SM私有的，L2是所有SM共享的。我习惯把L1看作「工作组内的黑板」，L2看作「全局公告栏」。
| 特性 | CPU缓存 | GPU缓存 |
|---|---|---|
| 设计目标 | 降低单线程延迟 | 提高整体吞吐量 |
| 一致性模型 | 强一致性（MESI） | 弱一致性（需显式同步） |
| 容量 | L1 32KB，L3几十MB | L1 128KB，L2几MB |
| 延迟 | 几个时钟周期 | 几十到几百周期 |
| 预取 | 硬件自动预取 | 主要靠程序员手动管理 |
这里有个关键点：GPU的L1缓存和Shared Memory在物理上是同一块硬件。你可以把它配置成纯缓存，也可以配置成程序员可控的Shared Memory。我个人习惯在需要频繁复用的数据上用Shared Memory，在随机访问的场景下用缓存模式。
2.3 为什么GPU需要缓存？
你可能会问：GPU不是靠「隐藏延迟」吃饭的吗？既然它能通过切换线程来掩盖内存访问延迟，为什么还需要缓存？
好问题。我来解释一下。
GPU确实可以通过「零开销线程切换」来隐藏延迟。但隐藏不等于消除。当所有线程都在等内存时，GPU就「卡住」了。缓存的作用就是减少这种「全员等待」的情况。
GPU缓存的三大作用：
- 减少全局内存访问：L2缓存命中一次，能省下几百个周期的延迟
- 支持数据复用：同一个warp内的线程访问相邻地址时，缓存能合并访问
- 降低带宽压力：缓存命中率高，全局内存的带宽需求就小
我记得有一次优化一个矩阵乘法的kernel。原始版本直接访问全局内存，性能只有50 GFLOPS。后来把分块数据放到Shared Memory（也就是L1缓存），性能飙到了800 GFLOPS。你看，缓存用好了，效果立竿见影。
2.4 一致性模型的差异
CPU和GPU在缓存一致性上走的是完全不同的路。
CPU采用强一致性模型。一个核改了数据，其他核立刻能看见。这靠的是MESI协议和总线嗅探机制。代价呢？核越多，一致性开销越大。所以CPU最多也就几十个核。
GPU采用弱一致性模型。一个线程改了数据，其他线程不一定能马上看见。你需要显式地调用__syncthreads()或者__threadfence()来保证可见性。这听起来麻烦，但好处是扩展性好——GPU可以轻松塞进几千个核心。
// GPU上的显式同步示例
__global__ void example_kernel(float* data) {
__shared__ float shared_data[256];
int tid = threadIdx.x;
shared_data[tid] = data[tid];
// 必须同步，否则其他线程可能读到未写入的值
__syncthreads();
// 现在可以安全地读取其他线程写入的数据
float val = shared_data[(tid + 1) % 256];
data[tid] = val;
}避坑指南：我曾经在GPU代码中忘记加__syncthreads()，结果数据时对时错。调试了两天才发现是同步问题。记住：GPU不会自动帮你保证缓存一致性，这是程序员的责任。
2.5 实际项目中的选择策略
说了这么多理论，咱们来点实际的。在项目中怎么选？
我一般遵循这几个原则：
- 数据只读且访问模式规则：用纹理缓存或常量缓存，硬件会自动优化
- 数据被反复重用：手动管理Shared Memory，比自动缓存更可控
- 数据访问随机且无规律：用L1/L2缓存，让硬件去处理
- 数据只在单个warp内使用：用寄存器，最快
嗯，说白了就是：能放寄存器的放寄存器，能放Shared Memory的放Shared Memory，剩下的交给缓存。这个优先级顺序，我建议你记下来。
最后说一句：GPU缓存不是万能的。它解决的是「数据局部性」问题，解决不了「算法复杂度」问题。如果你的算法本身就不合理，再好的缓存也救不了。这个道理，我在无数项目中验证过。