2、GPU内存层次结构：全局内存、共享内存、寄存器、常量内存、纹理内存的对比
说到GPU编程，内存管理永远是绕不开的核心话题。我经常跟团队里的新人说：你写CUDA代码的水平，很大程度上取决于你对内存层次的理解深度。为什么？因为数据在GPU里怎么流动、存在哪里，直接决定了你的程序能跑多快。
这一章，咱们就把GPU的五种主要内存掰开揉碎了讲清楚。我会结合自己踩过的坑，帮你建立起清晰的内存使用直觉。
2.1 一张图看懂全局架构
先看下面这张图，它概括了GPU内存的完整层次结构。我建议你把它当作「地图」来用——写代码时随时对照，就知道数据该放哪了。
从这张图你能看到：越靠近计算单元（SM）的内存，速度越快，但容量越小。这是个经典的「容量-速度」权衡问题。我刚开始做CUDA时，总想把所有数据都塞进共享内存，结果发现根本装不下——这就是没理解层次结构的代价。
2.2 五种内存逐一拆解
2.2.1 全局内存——最常用，也最慢
全局内存是GPU里的大仓库。所有线程都能读写，容量通常有8GB、16GB甚至更大。但代价是延迟高得吓人——一次访问要400到800个时钟周期。
我个人习惯把全局内存比作「硬盘」：能存很多东西，但每次读写都像在等快递。你想想看，如果你的核心循环里频繁访问全局内存，那性能基本就交代了。
关键点：全局内存访问必须满足「合并访问」条件才能发挥带宽。简单说就是：相邻线程访问相邻地址。我见过太多人因为没注意这个，导致带宽利用率只有10%。
我的经验：能用共享内存做数据复用的地方，就别反复读全局内存。我曾经优化过一个矩阵乘法，仅仅是把数据先搬进共享内存，性能就提升了8倍。
2.2.2 共享内存——程序员可控的缓存
共享内存是每个线程块私有的高速存储区。延迟只有30个周期左右，比全局内存快了一个数量级。但它有两个限制：容量小（通常48KB到164KB），而且线程块内共享。
共享内存最经典的用法就是做数据复用。比如做卷积运算时，每个线程需要访问输入数据的某个窗口。如果每个线程都去全局内存读，那就太蠢了。正确的做法是：先把整个窗口的数据加载到共享内存，然后所有线程从共享内存读。
// 共享内存典型用法：数据分块加载
__global__ void kernel(float* input, float* output, int N) {
__shared__ float s_data[BLOCK_SIZE]; // 声明共享内存
int tid = threadIdx.x;
int idx = blockIdx.x * blockDim.x + tid;
// 从全局内存加载到共享内存
if (idx < N) {
s_data[tid] = input[idx];
}
__syncthreads(); // 同步！确保所有线程加载完成
// 从共享内存读取，进行计算
// ... 这里做你的核心计算 ...
// 写回全局内存
if (idx < N) {
output[idx] = s_data[tid] * 2.0f;
}
}注意：使用共享内存时，__syncthreads() 是必须的。我曾经因为忘记加这个同步，导致线程A读到的数据还是上一轮计算的结果——debug了整整两天才发现。
2.2.3 寄存器——最快，但最稀缺
寄存器是GPU里最快的存储，访问延迟为0个周期。每个线程都有自己的私有寄存器，但数量非常有限——每个线程最多255个寄存器（不同架构略有差异）。
你可能会想：「寄存器这么快，那我多用点不就好了？」嗯，这里要注意：寄存器用多了，会降低SM上同时运行的线程数量。比如一个SM最多能跑2048个线程，但如果每个线程用了32个寄存器，那SM就只能跑1024个线程了。这叫「寄存器压力」。
我记得有一次优化一个深度学习算子，为了省寄存器，我把一些中间结果拆成多次计算。虽然计算量增加了，但因为SM上能同时跑的线程数翻倍了，最终性能反而提升了30%。这就是权衡的艺术。
2.2.4 常量内存——只读，但有缓存
常量内存是全局内存的一个特殊区域，大小只有64KB。它有两个特点：只读，而且有专用的缓存。当所有线程都读取同一个地址时，常量内存的效率极高——一次广播就能搞定。
常量内存最适合放那些「所有线程都用，但不会变」的数据。比如神经网络的权重、物理模拟中的常数参数等。
避坑指南：我曾经把一组查找表放在常量内存里，结果发现性能还不如全局内存。后来一查才知道：常量内存的缓存只有一次广播能力。如果不同线程读不同地址，那就退化成串行访问了。所以，常量内存只适合「所有线程读同一个地址」的场景。
2.2.5 纹理内存——为空间局部性而生
纹理内存最初是为图形渲染设计的，但在通用计算里也很有用。它也是只读的，但有专门的缓存机制，特别适合二维或三维的空间局部性访问。
什么叫空间局部性？就是相邻的线程访问相邻的地址。比如图像处理中，每个像素需要读取它周围3x3邻域的数据。这种情况下，纹理内存的缓存命中率非常高。
纹理内存还有一个绝活：硬件插值。如果你需要做双线性插值或三线性插值，纹理内存可以直接帮你算好，省掉一堆计算指令。
2.3 五种内存的对比表格
下面这个表格是我自己整理的，写代码时经常拿出来对照。你最好也记住这些关键数字。
| 内存类型 | 作用域 | 生命周期 | 读写权限 | 延迟 | 典型容量 | 缓存 |
|---|---|---|---|---|---|---|
| 全局内存 | 所有线程 | 程序运行期间 | 读/写 | 400-800 cycles | GB级 | L1/L2 |
| 共享内存 | 线程块内 | 内核执行期间 | 读/写 | ~30 cycles | 48-164 KB | 无 |
| 寄存器 | 单个线程 | 内核执行期间 | 读/写 | 0 cycles | 256 KB/block | 无 |
| 常量内存 | 所有线程 | 程序运行期间 | 只读 | ~10 cycles(缓存命中) | 64 KB | 专用缓存 |
| 纹理内存 | 所有线程 | 程序运行期间 | 只读 | ~100 cycles(缓存命中) | 取决于分配 | 专用缓存 |
2.4 实战中的选择策略
说了这么多理论，到底怎么选？我总结了一个简单的决策流程：
- 数据是否只读？ → 是 → 考虑常量内存或纹理内存
- 数据是否被线程块内多个线程复用？ → 是 → 用共享内存
- 数据是否每个线程私有且访问频繁？ → 是 → 尽量用寄存器
- 以上都不满足？ → 用全局内存，但注意合并访问
这个流程不是绝对的，但能帮你快速定位。我刚开始做优化时，就是靠这个流程一步步排查，很少走弯路。
一个小技巧：用 nvcc --ptxas-options=-v 编译时，会输出每个内核的寄存器使用量和共享内存使用量。我每次写完内核都会跑一下这个命令，看看资源使用是否合理。
2.5 一个综合示例
最后，给你看一个综合使用多种内存的代码片段。这是一个简化版的图像卷积核：
__constant__ float c_kernel[9]; // 卷积核放在常量内存
__global__ void convolution_2d(float* input, float* output, int width, int height) {
__shared__ float s_tile[TILE_SIZE + 2][TILE_SIZE + 2]; // 共享内存存图块
int tx = threadIdx.x;
int ty = threadIdx.y;
int gx = blockIdx.x * TILE_SIZE + tx;
int gy = blockIdx.y * TILE_SIZE + ty;
// 加载数据到共享内存（包含边界填充）
// ... 这里省略边界处理逻辑 ...
s_tile[ty + 1][tx + 1] = input[gy * width + gx];
__syncthreads();
// 计算卷积（从共享内存读，用常量内存的卷积核）
if (gx < width && gy < height) {
float sum = 0.0f;
for (int i = 0; i < 3; i++) {
for (int j = 0; j < 3; j++) {
sum += s_tile[ty + i][tx + j] * c_kernel[i * 3 + j];
}
}
output[gy * width + gx] = sum;
}
}这个例子中：卷积核放常量内存（所有线程读相同数据），图像块放共享内存（线程块内复用），中间变量用寄存器（编译器自动处理）。三种内存各司其职，性能自然就上去了。
好了，关于GPU内存层次结构就讲到这里。记住一句话：没有最好的内存，只有最合适的用法。下一章咱们会深入共享内存的bank conflict问题——那才是真正考验功力的地方。