内存层次结构与优化策略
大家好，我是你们的老朋友。今天我们来聊聊CUDA编程里最核心、也最考验功力的部分——内存层次结构。
说实话，我见过太多人写CUDA代码，核函数写得飞快，一跑起来性能惨不忍睹。为什么？十有八九是内存访问出了问题。GPU的计算能力再强，如果数据喂不进去，那也是白搭。这就像你开着一辆法拉利，结果堵在早高峰的二环上，再好的引擎也发挥不出来。
所以，理解GPU的内存体系，是每个CUDA程序员的必修课。我个人习惯把GPU内存比作一个「多级仓库」：离计算单元越近的，速度越快，但容量越小；离得越远的，容量大，但延迟高得吓人。咱们今天就把这个仓库的每个角落都翻一遍。
全局内存：最常用，也最坑
全局内存是GPU里容量最大的存储空间，所有线程都能访问。但它的延迟也是最高的——大概几百个时钟周期。你想想看，CPU访问L1缓存只要几个周期，这差距有多大。
全局内存最大的坑是什么？合并访问。
我在项目中遇到过这样一个案例：一个简单的向量加法，一开始写的版本性能只有理论峰值的10%。排查了半天，发现是内存访问模式不对。GPU一次从全局内存读取数据，是以128字节为单位的「事务」（transaction）。如果同一个warp里的32个线程，访问的地址是连续的、对齐的，那一次事务就能搞定。但如果大家东一榔头西一棒子，GPU就得拆成多个事务来读，带宽利用率直线下降。
合并访问的核心原则：
- 同一个warp的线程，访问连续的内存地址
- 起始地址要对齐到128字节
- 尽量使用
float4或int4等向量类型，一次读16字节
举个例子，假设我们要处理一个结构体数组（AoS）：
// 不推荐：AoS 布局，线程访问不连续
struct Particle {
float x, y, z, vx, vy, vz;
};
Particle particles[N];
// 线程 i 访问 particles[i].x —— 地址跳跃很大
// 推荐：SoA 布局，连续访问
struct Particles {
float *x, *y, *z;
float *vx, *vy, *vz;
};
// 线程 i 访问 x[i] —— 地址连续
说白了，把数据排布从AoS改成SoA，很多时候就能带来几倍的性能提升。我刚开始做CUDA优化时，就靠这一招，把一个分子动力学模拟从2小时缩短到了20分钟。
共享内存：程序员手中的「瑞士军刀」
共享内存是片上的SRAM，延迟极低（大概几个周期），容量只有几十KB。它是同一个block内所有线程共享的。说白了，这就是你手动控制的缓存。
共享内存最大的价值在于数据复用。比如矩阵乘法，每个输出元素需要读取一整行和一整列的数据。如果不做优化，每个线程都去全局内存读，带宽很快就撑爆了。正确的做法是：把数据分块加载到共享内存里，让block内的线程反复使用。
我的经验：共享内存的容量很宝贵，一般就48KB（有些卡可以配到96KB或164KB）。分配的时候要精打细算，别一股脑全占了。可以用 cudaFuncSetAttribute 调整每个block的共享内存大小。
共享内存还有一个大坑——Bank Conflict。
共享内存被分成了32个bank（对应warp的32个线程）。每个bank每个周期只能响应一个请求。如果同一个warp里多个线程访问同一个bank的不同地址，就会发生bank conflict，导致访问串行化。
举个例子：
// 假设 shared float s[32][32];
// 线程 i 访问 s[i][i] —— 所有线程访问不同的bank，无冲突
// 线程 i 访问 s[i][0] —— 所有线程访问bank 0，32路冲突！
// 线程 i 访问 s[i][1] —— 所有线程访问bank 1，32路冲突！
怎么避免？加padding。比如把 s[32][32] 声明成 s[32][33]，这样每行多了一个元素，列索引相同的元素就分散到了不同的bank里。
我记得有一次优化一个3D卷积的核函数，共享内存的bank conflict从16路降到了0路，性能直接翻倍。嗯，这种成就感，只有做过的人才能体会。
寄存器：最快，但别滥用
寄存器是GPU里最快的存储，没有之一。每个线程都有自己的寄存器文件，访问延迟为0（一个周期内完成）。
但寄存器数量是有限的。每个SM（流多处理器）的寄存器总数是固定的，比如每个SM有65536个寄存器。如果每个线程用32个寄存器，那一个SM最多能驻留65536/32=2048个线程。如果每个线程用64个寄存器，那只能驻留1024个线程。
注意：寄存器用太多，会降低occupancy（驻留线程数），反而可能拖慢性能。我见过有人为了省共享内存，把所有中间变量都塞进寄存器，结果occupancy掉到25%，性能还不如用共享内存的方案。
怎么控制寄存器用量？编译时加 -maxrregcount=32 参数，或者用 __launch_bounds__ 限定每个block的最大线程数。编译器会帮你做寄存器溢出（spill）到本地内存，但溢出后的性能会下降很多。
本地内存：寄存器的「备胎」
本地内存听起来像「本地」，其实它和全局内存是同一块物理存储，延迟一样高。只是每个线程有自己私有的地址空间。
什么时候会用本地内存？当寄存器不够用的时候，编译器会把一些局部变量「溢出」到本地内存。还有就是数组索引不是编译期常量时，也可能会被放到本地内存。
怎么避免？减少每个线程的寄存器使用量，或者把大数组放到共享内存里。我曾经调试过一个核函数，性能只有预期的30%，一看PTX代码，发现一个局部数组被放到了本地内存，每次访问都是几百周期的延迟。改成共享内存后，性能恢复正常。
常量内存：小而快，适合广播
常量内存是全局内存的一部分，但有专门的缓存（常量缓存），每个SM只有64KB。它的特点是：所有线程访问同一个地址时，速度极快（一个周期）。但如果不同线程访问不同地址，就会串行化。
所以常量内存最适合存什么？查找表、系数、参数这些所有线程都要读的「只读」数据。
举个例子，一个softmax核函数里的指数表，或者一个物理模拟里的物理常数，放在常量内存里就很合适。
// 声明常量内存
__constant__ float lookup_table[256];
// 在主机端初始化
cudaMemcpyToSymbol(lookup_table, host_table, sizeof(float) * 256);
// 核函数里直接访问
__global__ void kernel() {
float val = lookup_table[idx]; // 所有线程读同一个地址，快！
}
纹理内存：为空间局部性而生
纹理内存最初是为图形学设计的，但它的一些特性在通用计算里也很有用：
- 硬件插值：可以自动做双线性/三线性插值
- 边界处理：支持 clamp、wrap 等寻址模式
- 缓存优化：针对2D空间局部性做了优化
什么时候用纹理内存？当你的访问模式是非对齐的、随机分布的，比如图像处理、体渲染、光线追踪。纹理缓存的2D局部性比全局内存的1D合并访问更适合这类场景。
我个人习惯：能用全局内存合并访问解决的，就不用纹理内存。因为纹理内存的读取有额外开销（地址转换、格式转换）。但遇到不规则访问时，纹理内存往往是救命稻草。
一张图看懂内存层次
下面这张SVG图，是我自己总结的内存层次结构。你可以把它当作「优化路线图」：
优化策略总结
好了，讲了这么多，我给大家总结几条实战经验：
- 先看全局内存：确保合并访问。这是性价比最高的优化，往往能带来数量级的提升。
- 再看共享内存：如果有数据复用，果断用共享内存。注意bank conflict，加padding解决。
- 控制寄存器用量：别让occupancy掉得太低。一般建议每个线程不超过32个寄存器。
- 常量内存和纹理内存：特定场景下使用。常量内存适合广播，纹理内存适合不规则2D访问。
- 本地内存是敌人：如果发现核函数性能异常低，检查PTX代码，看有没有变量被溢出到本地内存。
我曾经踩过的坑：有一次优化一个流体模拟的核函数，怎么调都调不好。后来用 nvcc --ptxas-options=-v 一看，发现每个线程用了80个寄存器，occupancy只有12.5%。我硬着头皮重构了代码，把一些中间结果放到共享内存里，寄存器降到32个，occupancy升到50%，性能提升了3倍。所以，别迷信「所有东西放寄存器最快」，有时候「慢一点但能并行更多」才是正解。
嗯，内存优化这块，说白了就是「权衡」二字。速度、容量、并行度，三者不可兼得。你需要根据具体的算法和数据访问模式，找到那个最合适的平衡点。多写、多测、多分析，慢慢就有感觉了。
想要更多资料，联系可可