显存类型详解：全局内存、共享内存、寄存器、本地内存、常量内存、纹理内存
各位同学，咱们今天来聊聊GPU里的“内存家族”。说实话，很多初学者写CUDA程序，上来就是cudaMalloc一把梭，全往全局内存里塞。这其实也没错，但性能嘛……嗯，你懂的。我刚开始做GPU加速那会儿，也踩过这个坑。后来被项目里的性能瓶颈逼着，才把这些内存类型一个个啃透。
今天我就带你把这六种内存类型捋一遍。搞懂了它们，你的核函数性能至少能翻个倍。
1. 全局内存（Global Memory）—— 最熟悉，也最慢
全局内存是GPU里容量最大的显存，所有线程都能访问。你平时用cudaMalloc分配的就是它。但它的延迟很高，大概400-800个时钟周期。
核心特点：
- 容量最大（几个GB到几十GB）
- 所有线程可读写
- 延迟最高，需要合并访问才能发挥带宽
- 没有缓存（L2缓存除外，但那是硬件自动管理的）
我个人习惯，全局内存只用来放输入数据和最终结果。中间计算能不用就不用。为什么？因为慢啊！你想想看，一个线程读一次全局内存的时间，够寄存器算几十次加法了。
避坑指南：
我曾经在项目里把一个大数组反复读写全局内存，结果性能惨不忍睹。后来改成先用共享内存做中间缓存，性能直接提升了5倍。记住：全局内存是“慢车道”，能少走就少走。
2. 共享内存（Shared Memory）—— 线程块的“私房钱”
共享内存是每个线程块私有的，块内所有线程都能访问。它的速度比全局内存快得多，延迟只有几十个时钟周期。
| 特性 | 说明 |
|---|---|
| 容量 | 每个SM通常48KB-164KB（取决于架构） |
| 延迟 | 约20-30个时钟周期 |
| 作用域 | 仅限同一个线程块 |
| 生命周期 | 线程块执行期间 |
共享内存最经典的用法就是做数据复用。比如矩阵乘法，你把子矩阵加载到共享内存里，然后反复用它做计算，避免反复读全局内存。
// 声明共享内存
__shared__ float s_data[32][32];
// 从全局内存加载到共享内存
s_data[tx][ty] = g_data[tx + blockIdx.x * blockDim.x][ty + blockIdx.y * blockDim.y];
__syncthreads(); // 同步！别忘了
// 然后就可以快速访问s_data了
float val = s_data[i][j]; // 快得很注意：共享内存用完后一定要加__syncthreads()同步。我见过太多人因为忘记同步，读到了脏数据，调试到怀疑人生。
3. 寄存器（Register）—— 最快，但最稀缺
寄存器是GPU里最快的存储，没有之一。每个线程都有自己的寄存器，延迟为0（一个时钟周期内完成）。但它的数量非常有限，每个线程通常只有32-255个（取决于架构和编译器设置）。
说白了，寄存器就是线程的“手边工具”。你定义的局部变量，只要编译器觉得合适，就会放到寄存器里。
关键点：
- 速度最快，零延迟
- 每个线程私有
- 数量有限，用多了会“溢出”到本地内存
- 编译器自动管理，但你可以用launch_bounds_提示
我在项目中遇到过一个问题：核函数里定义了太多局部变量，结果寄存器不够用，编译器悄悄把它们“溢出”到了本地内存。性能直接掉了30%。后来我精简了变量，或者用launch_bounds_限制线程数，才把寄存器占用压下来。
4. 本地内存（Local Memory）—— 寄存器的“替补”
本地内存听起来像“本地”的，其实它物理上就在全局内存里。只是每个线程私有的。当寄存器不够用时，编译器会把一些变量放到本地内存里。
为什么会这样？因为寄存器是有限的。你一个线程用了太多变量，寄存器放不下，就只能放到慢速的本地内存里。说白了，本地内存就是“寄存器的溢出区”。
避坑指南：
我曾经写过一个核函数，里面定义了一个大数组：float temp[256]。结果编译器直接把它放到了本地内存里，每次访问都慢如龟速。后来我改成用共享内存，性能才恢复正常。记住：大数组、或者编译器无法确定索引的数组，很容易被放到本地内存。
5. 常量内存（Constant Memory）—— 只读，但带缓存
常量内存是只读的，所有线程都能访问。它最大的特点是带有缓存，而且缓存是广播式的——如果所有线程读同一个地址，那速度极快。
常量内存容量很小，通常只有64KB。适合放一些不会变的参数，比如滤波器系数、查找表之类的。
// 声明常量内存
__constant__ float c_filter[64];
// 在主机端初始化
cudaMemcpyToSymbol(c_filter, host_filter, sizeof(float) * 64);
// 核函数里直接读
float val = c_filter[idx]; // 如果所有线程读同一个idx，那快得飞起我个人习惯，只要数据是只读的、且所有线程访问同一个地址，我就用常量内存。比如做图像处理时的卷积核，放在常量内存里再合适不过了。
6. 纹理内存（Texture Memory）—— 为空间局部性而生
纹理内存最初是为图形学设计的，但用在通用计算里也很香。它也是只读的，但它的缓存优化是针对空间局部性的——也就是说，如果你访问的地址在空间上相邻，缓存命中率会很高。
纹理内存还有一些特殊功能，比如硬件插值、边界处理等。做图像处理、科学计算时特别好用。
| 特性 | 说明 |
|---|---|
| 只读 | 所有线程只能读，不能写 |
| 缓存 | 针对2D/3D空间局部性优化 |
| 特殊功能 | 硬件插值、边界钳制、归一化坐标 |
| 适用场景 | 图像处理、随机访问模式、需要插值的计算 |
实战经验：
我记得有一次做医学图像重建，数据访问模式完全是随机的。用全局内存读，缓存命中率低得可怜。后来换成纹理内存，利用它的空间局部性缓存，性能提升了2倍多。嗯，纹理内存在这种场景下就是神器。
总结一下
好了，六种内存类型都讲完了。我帮你理一下思路：
- 寄存器：最快，但少。局部变量优先放这里。
- 共享内存：块内共享，适合数据复用。手动管理，要加同步。
- 全局内存：容量大，速度慢。只放输入输出。
- 本地内存：寄存器的替补，尽量别用。
- 常量内存：只读，带广播缓存。适合常量参数。
- 纹理内存：只读，空间局部性优化。适合随机访问和图像处理。
你想想看，如果你能把数据合理地分配到这些内存里，性能怎么可能差？我每次写核函数前，都会先画个数据流图，标清楚哪些数据放哪里。这个习惯帮我省了不少调试时间。
最后一个小技巧：用nvcc的--ptxas-options=-v选项编译，可以看到每个核函数用了多少寄存器、共享内存。如果发现寄存器用多了，可以适当减少线程数或者精简变量。我曾经靠这个选项，把一个核函数的寄存器占用从64降到了32，性能直接翻倍。
下一章咱们聊聊内存访问模式，特别是“合并访问”这个坑。嗯，到时候你就知道为什么全局内存那么慢了。