CUDA 内存层次：五大内存的实战理解

咱们做 GPU 编程，说白了就是在跟内存打交道。我刚转 CUDA 那会，以为把数据扔到显存里就完事了，结果性能一塌糊涂。后来才明白，搞懂内存层次，比学什么并行算法都重要。

我先帮你梳理一张总图，你大概看一眼五个家伙在芯片上的位置关系，然后再一个一个拆开讲。

1. 寄存器：最快的，也是最「小气」的

寄存器就在 SM 内部，跟计算单元贴在一起。访问延迟只有 1 个周期，比其它内存快一两个数量级。但我告诉你，这玩意儿特别「小气」——每个线程能分到的数量是有限的。

很多新手上来就声明一大堆局部变量，结果把寄存器用爆了。一旦爆了，数据会被「spill」到全局内存，性能直接崩盘。**我个人习惯在 kernel 开头加 --maxrregcount 编译选项卡一下寄存器上限，宁可让编译器多调度几次，也比 spill 强。**

2. 共享内存：手动管理的片上缓存

共享内存也在芯片上，所以延迟很低，5 到 30 个周期。但它需要你用 __shared__ 手动声明，完事还得自己去同步。

我记得第一次写矩阵分块乘法的 kernel，靠共享内存做 tile 缓存，速度比直接读写全局内存快了将近 10 倍。**但千万别忽略 bank conflict**——同一 warp 里多个线程同时访问同一个 bank 的不同地址，本来并行的事就串行了，性能掉得很快。

**我以前在项目里调优一个规约 kernel，死活差 30% 的性能没跑满。后来用 cuda-memcheck 一查，发现是 bank conflict 打满了，改成无冲突的 stride + padding 布局才解决问题。**

3. 全局内存：仓库大，但离得远

全局内存就是显存，容量最大（GB 级），但延迟 400-800 个周期。你想想看，CPU 这边 DRAM 也就 100 纳秒左右，GPU 全局内存差不多也是这个量级，但因为 GPU 线程多，总带宽反而非常恐怖，HBM2e 能到 TB/s 级别。

用好全局内存的核心，只有两个字：**合并**。你得保证同一个 warp 里 32 个线程访问的地址是连续的，这样才能一次搞定一次事务。我记得有个学员把一个 2D 数组按行写入，性能惨不忍睹；改成行优先存储 + 连续线程索引，直接提速 4 倍。

**我自己的经验：永远先去检查访存模式。用 ncu --set full 看 L1/TEX cache 命中率，如果低于 60%，多半是合并出了问题。**

4. 常量内存：只读，但有广播优势

常量内存只有 64 KB，但带了一个专门的缓存。最重要的特点是：**warp 里所有线程访问同一个地址时，只发一次请求，然后广播给所有线程。**

这个特性特别适合存滤波器系数、查找表、神经网络权重这种东西。**我在做音频处理的 FIR 滤波器时，把 64 个系数丢到常量内存里，比全局内存快了一倍还多。**

但要注意，如果你用常量内存存了太多不同的值，每个线程都去读不同的地址，那缓存就炸了，退化成串行访问。**我建议只放「全 warp 一致」的只读数据，否则不如用全局内存配合只读纹理。**

5. 纹理内存：为空间局部性而生

纹理内存这个名称是从图形学来的，但它对非图形计算同样有用。它专门优化了 2D / 3D 的空间局部性访问——你访问点 (x, y)，它会把邻居也预取到缓存里。

我当时做流场可视化，需要随机地在 2D 网格上查样，全局内存的 cache miss 高得离谱。换成纹理内存后，因为 tex cache 天然支持 2D 空间局部性，缓存命中率直接从 30% 涨到 80%。**而且纹理内存还自带硬件插值和边界处理，省了我手写 clamp 的功夫。**

不过实话实说，纹理内存只适合「读多写少 + 随机空间访问」的模式。如果你的访问模式是线性的、连续的，全局内存合并访问反而比纹理快。

选型对照表

避坑指南

• 别一股脑把所有变量都扔共享内存——只有被多个线程反复读写的才值得。我曾经见有人把每个线程的私有中间结果也丢共享内存，结果同步开销直接吃掉所有收益。
• 常量内存用不了大数组——64 KB 上限，超过请用 const __restrict__ 全局指针，配合编译器优化。
• 纹理内存不适合常规数组遍历——连续访问不如全局内存合并，且写纹理内存要走 surface API，稍微麻烦一点。
• 寄存器 spill 是性能杀手——如果你发现 kernel 里大量变量后面带 st.local / ld.local 指令，说明寄存器爆了，赶紧减小线程块尺寸或拆 kernel。

一句话总结

你写 CUDA 时，脑子里始终有一张金字塔图：寄存器是手边抽屉，共享内存是办公桌，全局内存是仓库，常量内存是广播公告栏，纹理内存是带磁性的工作台。每次数据放到哪一层，问问自己「这个值被谁用、怎么用、用几次」，答案自然就出来了。