4、CUDA内存模型：全局内存、共享内存、寄存器、本地内存、常量内存、纹理内存
说到CUDA编程，内存模型是绕不开的核心话题。我刚开始学的时候，觉得这东西不就是显存嘛，有啥好讲的？结果第一次写核函数，性能惨不忍睹，才意识到——内存访问模式决定了你的GPU能跑多快。
说白了，GPU计算就是数据搬运。你把数据搬得快，计算就快；搬得慢，计算核心再强也白搭。今天咱们就把CUDA的六种内存掰开揉碎讲清楚。
4.1 全局内存（Global Memory）
全局内存是GPU里最大的存储空间，所有线程都能访问。我习惯叫它「显存」，虽然严格来说不完全等价。
关键特性：
- 容量最大（通常几GB到几十GB）
- 延迟最高（几百个时钟周期）
- 所有线程可读写
- 主机端可通过cudaMemcpy访问
我在项目中遇到过一个问题：数据量太大，全局内存不够用。那时候只能分块处理，每次搬运一部分到共享内存里算。嗯，这里要注意——全局内存的访问一定要合并。
什么叫合并访问？就是相邻线程访问相邻地址。举个例子：
// 好的访问模式：合并访问
float value = data[threadIdx.x]; // 线程0访问data[0]，线程1访问data[1]...
// 坏的访问模式：非合并访问
float value = data[threadIdx.x * 1024]; // 线程0访问data[0]，线程1访问data[1024]...
我曾经因为没注意这个，一个简单的向量加法慢了10倍。你想想看，几百个线程同时访问内存，如果地址不连续，内存控制器得累死。
4.2 共享内存（Shared Memory）
共享内存是我最喜欢的部分。它就在每个SM（流多处理器）内部，延迟极低（大概1-2个时钟周期）。
使用技巧：
- 容量有限（通常48KB或96KB）
- 同一个block内的线程共享
- 需要手动管理数据搬运
- 注意bank conflict问题
我建议你把共享内存当成「程序员手动控制的缓存」。全局内存太慢，寄存器太少，共享内存正好在中间。
来看一个矩阵转置的例子：
__global__ void transpose_v1(float *in, float *out, int width) {
__shared__ float tile[TILE_SIZE][TILE_SIZE];
int x = blockIdx.x * TILE_SIZE + threadIdx.x;
int y = blockIdx.y * TILE_SIZE + threadIdx.y;
// 加载到共享内存
tile[threadIdx.y][threadIdx.x] = in[y * width + x];
__syncthreads();
// 转置后写回
out[x * width + y] = tile[threadIdx.x][threadIdx.y];
}
这里有个坑：__syncthreads() 必须加。我曾经漏掉它，结果数据还没加载完就开始计算，出来的结果全是乱的。调试了一整天才发现。
4.3 寄存器（Register）
寄存器是GPU里最快的存储，没有之一。每个线程私有的，编译器会自动分配。
注意事项：
- 每个线程最多255个寄存器（不同架构有差异）
- 寄存器溢出会使用本地内存（变慢）
- 编译器会尽量优化寄存器使用
我记得有一次优化一个卷积核，发现性能上不去。用 --ptxas-options=-v 一看，每个线程用了80多个寄存器，导致占用率很低。后来手动拆分代码，减少临时变量，把寄存器压到32个，性能直接翻倍。
你想想看，寄存器越多，SM能同时运行的线程就越少。这就是占用率（occupancy）的概念。
4.4 本地内存（Local Memory）
本地内存听起来像本地存储，其实它物理上在全局内存里。只是逻辑上每个线程私有。
什么时候会用本地内存？
- 寄存器不够用的时候（溢出）
- 访问大型数组（编译器自动判断）
- 某些复杂的索引计算
避坑指南：我曾经写了一个核函数，里面定义了一个很大的局部数组 float temp[1024]。结果编译器把它放到了本地内存，每次访问都要走全局内存的延迟，性能惨不忍睹。后来改成共享内存，速度提升了5倍。
怎么避免？尽量控制局部变量的大小，能用共享内存就用共享内存。
4.5 常量内存（Constant Memory）
常量内存是个特殊的存在。它容量小（64KB），但有专门的缓存，访问速度很快。
使用场景：
- 存储不会改变的参数（如卷积核权重）
- 所有线程访问相同的数据
- 数据量小于64KB
__constant__ float filter[256];
// 主机端初始化
cudaMemcpyToSymbol(filter, host_filter, sizeof(float) * 256);
__global__ void apply_filter(float *data) {
// 所有线程访问同一个filter
float val = data[threadIdx.x] * filter[threadIdx.x % 256];
}
我习惯把模型推理时的固定参数放在常量内存里。比如归一化的均值和方差，这些值不会变，而且所有线程都要用。
4.6 纹理内存（Texture Memory）
纹理内存最初是为图形处理设计的，但在通用计算里也有妙用。
它的特点：
- 有硬件插值功能
- 对二维/三维空间访问友好
- 有专门的缓存
- 适合随机访问模式
个人经验：我在做图像处理时，经常用纹理内存做双线性插值。硬件直接帮你算好了，比自己写插值函数快得多。不过要注意，纹理内存的写入很麻烦，一般只读。
4.7 各内存对比总结
| 内存类型 | 位置 | 容量 | 延迟 | 作用域 | 缓存 |
|---|---|---|---|---|---|
| 全局内存 | 显存 | GB级 | 高 | 所有线程 | L2 |
| 共享内存 | SM内部 | KB级 | 低 | Block内 | 无 |
| 寄存器 | SM内部 | 每个线程有限 | 极低 | 线程私有 | 无 |
| 本地内存 | 显存 | 取决于溢出 | 高 | 线程私有 | L1/L2 |
| 常量内存 | 显存 | 64KB | 低（有缓存） | 所有线程 | 专用缓存 |
| 纹理内存 | 显存 | GB级 | 中 | 所有线程 | 专用缓存 |
4.8 实战建议
说了这么多，到底怎么选？我总结几条经验：
- 优先用寄存器：能放寄存器就别放别处
- 共享内存做缓存：数据复用率高时，手动搬进共享内存
- 常量内存存参数：所有线程读相同数据时用
- 纹理内存做插值：图像处理、空间查询时考虑
- 避免本地内存：编译器提示寄存器溢出时，优化代码结构
最后提醒：别迷信理论延迟。实际性能要看内存带宽利用率。我曾经优化一个程序，把所有数据都放共享内存，结果因为bank conflict反而更慢了。所以，先分析访问模式，再选内存类型。
嗯，内存模型就讲到这里。下一章咱们聊聊线程同步与原子操作，这可是多线程编程的难点。到时候我会分享一个让我熬夜到凌晨3点的bug案例，保证让你印象深刻。