第四章：内存层次优化——全局内存、共享内存、寄存器、本地内存、常量内存、纹理内存
说到GPU性能优化，内存这块儿是绕不开的坎儿。我见过太多人，写出来的CUDA代码在逻辑上完全正确，但跑起来就是慢得让人抓狂。问题出在哪儿？十有八九是内存访问没整明白。
GPU的内存不是铁板一块，它分了好几个层次。每个层次的速度、容量、用途都不一样。说白了，这就是个金字塔结构——越靠近计算单元的内存越快，但容量也越小。咱们今天就把这六个层次掰开揉碎了讲清楚。
4.1 全局内存——最基础也最慢
全局内存是GPU里容量最大的存储空间，所有线程都能访问。但代价就是慢，真的慢。一次全局内存访问大概要400到800个时钟周期，这期间计算单元只能干等着。
我刚开始做优化的时候，总觉得把数据放到全局内存里就完事了。后来发现，性能瓶颈全在这儿。你想想看，如果每个线程都随机访问全局内存，那效率简直惨不忍睹。
关键点：全局内存访问要满足合并访问（coalesced access）原则。也就是说，同一个warp里的32个线程，最好访问连续的内存地址。这样硬件可以把多次访问合并成一次大的内存事务。
举个例子，假设你要处理一个二维数组：
// ❌ 糟糕的访问模式——按列访问
for (int i = 0; i < N; i++) {
float val = data[i * width + col]; // 每个线程访问不同行
}
// ✅ 好的访问模式——按行访问
float val = data[row * width + threadIdx.x]; // 连续线程访问连续地址我在项目中遇到过有人写了个矩阵转置，性能差得离谱。一查，就是全局内存访问没合并。改成按行读取、按列写入后，速度直接翻了4倍。
4.2 共享内存——线程块的"小团队缓存"
共享内存是每个线程块私有的高速缓存，容量通常只有几十KB。但它快啊，延迟只有30个时钟周期左右，比全局内存快了一个数量级。
共享内存最常用的场景就是数据复用。比如做矩阵乘法时，每个子矩阵会被多次读取。把它加载到共享内存里，就能避免反复访问全局内存。
我的习惯：写kernel之前，先想想哪些数据会被多次使用。如果有，就考虑用共享内存做缓存。但要注意，共享内存容量有限，别塞太多东西导致occupancy下降。
共享内存还有个坑——bank conflict。共享内存被分成32个bank，如果多个线程同时访问同一个bank的不同地址，就会发生冲突，导致访问串行化。
// 共享内存声明
__shared__ float s_data[256];
// 无bank conflict的访问——每个线程访问不同bank
float val = s_data[threadIdx.x]; // 没问题
// 有bank conflict的访问——线程0和线程16访问同一个bank
float val = s_data[threadIdx.x * 2]; // 步长为2，可能冲突我曾经调试过一个图像滤波的kernel，性能一直上不去。后来发现是共享内存的bank conflict太严重。调整了数据布局后，性能提升了30%。
4.3 寄存器——最快但最稀缺
寄存器是离计算单元最近的内存，访问延迟几乎为零。每个SM有65536个寄存器（以现代GPU为例），但分配到每个线程就很少了。
你想想看，如果一个SM要跑2048个线程，那每个线程只能分到32个寄存器。一旦你的kernel用了超过这个数，多余的变量就会被"溢出"到本地内存——那速度就掉到全局内存的水平了。
注意：别为了省寄存器而把变量声明成数组。编译器有时候会把数组放到本地内存里，反而更慢。我建议用--ptxas-options=-v编译选项查看寄存器使用情况。
我个人习惯是，尽量把频繁使用的变量放在寄存器里。比如循环计数器、累加器这些。但也要控制变量数量，别让编译器被迫做寄存器溢出。
4.4 本地内存——寄存器的"替补队员"
本地内存听起来像是"本地"的，但实际上它存储在全局内存里，只是每个线程私有的。当寄存器不够用时，编译器会自动把一些变量放到本地内存里。
哪些变量容易进本地内存？大型数组、结构体、或者编译器无法确定访问模式的变量。说白了，就是那些寄存器装不下的东西。
我记得有一次优化一个复杂的kernel，发现性能比预期差很多。用nvcc --ptxas-options=-v一看，本地内存用了不少。后来通过拆分kernel、减少每个线程的变量数量，才把本地内存用量降下来。
避坑指南：我曾经把一个很大的局部数组声明在kernel里，结果每个线程都占用了大量本地内存。改成用共享内存后，性能提升明显。记住，本地内存不是"本地"的，它还是全局内存的速度。
4.5 常量内存——只读且带缓存
常量内存是只读的，容量只有64KB。但它有专门的缓存，当所有线程访问同一个地址时，效率极高。
常量内存最适合存那些不会改变的参数，比如卷积核的权重、归一化系数等。所有线程都读同一个值，缓存命中率接近100%。
// 声明常量内存
__constant__ float c_filter[1024];
// 在host端初始化
cudaMemcpyToSymbol(c_filter, host_filter, sizeof(float) * 1024);
// 在kernel中使用——所有线程读同一个地址
float val = c_filter[idx]; // 高效！但要注意，如果每个线程读不同的地址，常量内存就废了。因为常量内存的缓存是广播式的，只适合"所有人读同一个值"的场景。
4.6 纹理内存——为空间局部性而生
纹理内存最初是为图形学设计的，但在通用计算里也很有用。它利用空间局部性，对二维或三维数据的随机访问有优化。
纹理内存的好处是：
- 硬件自动处理边界情况（clamp、wrap等模式）
- 对非对齐访问有较好的性能
- 支持线性插值（对某些计算有用）
我一般在处理图像数据或者稀疏矩阵时用纹理内存。比如做图像卷积，用纹理内存读取像素，比全局内存快不少。
小技巧：纹理内存适合"读多写少"的场景。如果你的kernel主要是读取数据，而且访问模式不规则，试试纹理内存。但要注意，纹理内存是只读的。
4.7 实战建议——怎么选？
说了这么多，到底怎么用？我总结了几条经验：
| 场景 | 推荐内存 | 原因 |
|---|---|---|
| 数据量大，每个线程只读一次 | 全局内存 | 没必要搬来搬去，直接读就行 |
| 数据被多次复用 | 共享内存 | 减少全局内存访问次数 |
| 频繁使用的标量变量 | 寄存器 | 最快，但注意数量限制 |
| 所有线程读同一个值 | 常量内存 | 缓存命中率高 |
| 不规则访问的只读数据 | 纹理内存 | 空间局部性优化 |
嗯，内存优化这事儿，说白了就是"把数据放在离计算最近的地方"。但也不能盲目追求快，还得考虑容量和编程复杂度。我个人的经验是：先跑一个baseline，然后用profiler看看瓶颈在哪儿，再针对性地优化。
记住，没有银弹。每个场景都要具体分析。但只要你理解了这六个内存层次的特点，优化起来就有方向了。