2. CUDA内存架构全景
说实话，我刚接触CUDA那会儿，最头疼的就是内存模型。
全局内存、共享内存、寄存器、本地内存、常量内存、纹理内存……光名字就够绕的。更别提它们各自的延迟、带宽、作用域和生命周期了。我踩过不少坑，比如把频繁访问的数据放在全局内存里，结果性能惨不忍睹。后来才明白——搞懂存储层次，是GPU优化的第一课。
核心观点：GPU的存储层次，本质是用“距离计算单元越近，速度越快，容量越小”的物理规律，换取极致吞吐。你写的每一行代码，都在跟这个层次打交道。
2.1 一张图看懂GPU存储层次
先别急着背参数。我习惯把GPU存储想象成一个“金字塔”：
嗯，这张图我手绘过很多次。每次给团队新人讲优化，我都会先画这个金字塔。你想想看——越靠近顶端的存储，越快，但容量越小；越靠近底端，越慢，但容量越大。优化的本质，就是把数据尽量往顶端挪。
2.2 全局内存：GPU的“主存”
全局内存是GPU里容量最大的存储，也是延迟最高的。说白了，它就是显存（VRAM）。所有线程都能访问，但代价是每次访问要等几百个周期。
关键特性：
- 容量大：几GB到几十GB不等（取决于显卡）
- 延迟高：400~800个时钟周期
- 作用域：所有线程 + 主机端
- 生命周期：由程序员管理（cudaMalloc / cudaFree）
我曾经踩过的坑：刚学CUDA时，我写了一个矩阵乘法，所有数据都放在全局内存里。结果跑起来比CPU还慢。后来才发现——全局内存的带宽虽然高（比如RTX 4090有1TB/s），但如果不做合并访问（coalesced access），实际吞吐会掉到十分之一。
避坑指南：
- 尽量让相邻线程访问相邻地址（合并访问）
- 避免频繁的全局内存读写，能用共享内存就用共享内存
- 使用
__ldg()只读缓存加载（如果数据不变）
2.3 共享内存：线程块的“私房钱”
共享内存是我个人最喜欢的存储。它在芯片上，离计算单元很近，延迟只有几十个周期。同一个线程块里的所有线程都能访问它。
关键特性：
- 容量小：通常48KB~164KB per SM（取决于架构）
- 延迟低：约20~30个周期
- 作用域：同一个线程块内的所有线程
- 生命周期：线程块生命周期内
我的经验：共享内存是GPU优化的“杀手锏”。我在做卷积优化时，把输入特征图分块加载到共享内存里，然后让线程块内的所有线程反复使用这些数据。结果性能提升了5倍。说白了，就是用共享内存来“缓存”全局内存中的数据，减少对显存的访问。
使用示例：
// 声明共享内存（动态大小）
__shared__ float s_data[256];
// 或者静态声明
extern __shared__ float s_dynamic[];
// 从全局内存加载到共享内存
int tid = threadIdx.x;
s_data[tid] = global_data[blockIdx.x * blockDim.x + tid];
__syncthreads(); // 同步！确保所有线程加载完成
// 然后就可以快速访问共享内存了
float val = s_data[tid];提示：使用共享内存时，别忘了加 __syncthreads() 同步。我见过太多人因为忘记同步，导致数据竞争，结果算出来的结果时对时错。这种bug最难查。
2.4 寄存器：最快的存储，没有之一
寄存器是每个线程私有的，速度最快——一个周期就能访问。但数量极其有限。每个SM有成千上万个寄存器，但每个线程能用的数量是有限的（通常最多255个，但实际建议控制在32个以内）。
关键特性：
- 速度：~1个周期
- 容量：每个线程最多255个（但多了会 spill 到本地内存）
- 作用域：单个线程私有
- 生命周期：线程生命周期内
注意：如果你在kernel里声明了太多局部变量，寄存器不够用，编译器会把多余的变量“溢出”（spill）到本地内存（local memory）。本地内存其实在全局内存里，延迟很高。所以——寄存器不是越多越好。我习惯用 --ptxas-options=-v 编译选项查看寄存器使用情况，确保每个线程不超过32个寄存器。
2.5 常量内存与纹理内存：特殊用途的存储
这两个存储比较特殊，但用好了能带来惊喜。
常量内存：
- 只读，容量64KB
- 有专门的缓存（常量缓存）
- 适合存储系数、查找表等不变数据
- 所有线程访问同一地址时，速度极快（广播机制）
纹理内存：
- 只读，有专门的纹理缓存
- 支持硬件插值、边界处理
- 适合图像处理、非对齐访问
- 对二维/三维空间局部性友好
我的习惯：如果数据是只读的，而且所有线程访问同一个地址（比如滤波器系数），我会用常量内存。如果数据是只读的，但访问模式是随机的（比如图像像素），我会用纹理内存。这两种都能利用专门的缓存，减少全局内存压力。
2.6 缓存层次：L1、L2与只读缓存
现代GPU（从Kepler架构开始）有L1和L2缓存。L1是每个SM私有的，L2是所有SM共享的。
| 缓存类型 | 位置 | 容量 | 延迟 | 作用域 |
|---|---|---|---|---|
| L1缓存 | 每个SM内部 | 32KB~128KB（可配置） | ~30周期 | 单个SM |
| L2缓存 | 芯片全局 | 几MB（如RTX 4090有72MB） | ~200周期 | 所有SM |
| 只读缓存 | 每个SM内部 | 与L1共享 | ~30周期 | 单个SM |
你可能会问：L1缓存和共享内存是什么关系？嗯，在有些架构上（比如Turing、Ampere），L1和共享内存是物理上共享的同一块SRAM。你可以通过 cudaFuncSetAttribute 配置它们的大小比例。比如，你可以在编译时指定：
// 配置共享内存和L1的比例
cudaFuncSetAttribute(myKernel,
cudaFuncAttributeMaxDynamicSharedMemorySize,
49152); // 48KB共享内存核心建议：如果你的kernel需要大量共享内存，就多分给共享内存；如果不需要，就留给L1缓存。我一般会先默认配置跑一次，然后用profiler看L1命中率，再调整比例。没有银弹，只有实测。
2.7 总结：如何选择存储层次？
好了，说了这么多，到底怎么选？我总结了一个简单的决策流程：
- 数据是否只读？ → 是 → 用常量内存或纹理内存（看访问模式）
- 数据是否被线程块内多个线程复用？ → 是 → 用共享内存
- 数据是否每个线程私有且频繁访问？ → 是 → 用寄存器（注意不要溢出）
- 以上都不是？ → 用全局内存，但注意合并访问
说白了，优化的本质就是：把数据放在离计算最近的地方。我在项目中见过太多人，明明可以用共享内存，却非要从全局内存反复读取。结果就是——GPU利用率上不去，性能瓶颈在显存带宽。
记住一句话：全局内存是给“懒人”用的，共享内存是给“聪明人”用的，寄存器是给“高手”用的。你想想看，你想做哪一种？
最后一个小技巧：用 nvcc --ptxas-options=-v 编译时，会输出每个kernel的寄存器使用量、共享内存使用量、本地内存溢出量。我每次写完kernel都会看一眼这些数字，心里有个底。如果发现本地内存溢出（local memory spill），就说明寄存器用太多了，需要优化代码。
嗯，这一章就到这里。记住这张金字塔图，后面的所有优化技巧，都离不开它。