4、GPU内存管理：全局内存、共享内存、寄存器、本地内存、常量内存、纹理内存、内存分配与释放
说到GPU内存管理，我脑子里立刻浮现出刚入行时踩过的那些坑。那时候我天真地以为，显存嘛，不就是个大仓库，往里扔数据就行了。结果呢？性能惨不忍睹，还经常莫名其妙地报错。后来我才明白，GPU的内存体系远比想象中复杂——它不是一个简单的存储池，而是一套层次分明、各有分工的生态系统。
今天我们就来拆解这个系统。我会把每种内存的特点、用法、以及我亲身经历过的坑，都一一讲清楚。
4.1 全局内存（Global Memory）
全局内存是GPU里最大的存储空间，所有线程都能访问。说白了，它就是显存本体。容量大，但延迟也高——大概400-800个时钟周期。
特点：
- 可读可写，所有线程可见
- 生命周期由主机端控制，直到显式释放
- 未加修饰时，默认是全局内存
访问模式：
全局内存的性能，很大程度上取决于访问模式。我建议你记住一个原则：合并访问。什么意思呢？就是让同一个warp里的32个线程，访问连续的内存地址。这样硬件可以把多次访问合并成一次事务，效率翻倍。
关键点：全局内存的延迟是GPU性能的主要瓶颈之一。优化访问模式，比优化计算量更有效。
我的经验：有一次我处理一个矩阵转置的kernel，性能死活上不去。后来发现是访问模式不对——按列读取时，线程间地址跨度太大，导致合并访问失效。改成按行读取后，性能提升了3倍。嗯，这个教训我记得很清楚。
4.2 共享内存（Shared Memory）
共享内存是GPU里的一块“小快灵”存储。它在每个SM内部，容量很小（通常几十KB），但速度极快——延迟只有几个时钟周期。
特点：
- 同一个block内的线程共享
- 速度接近寄存器
- 需要程序员手动管理数据搬入搬出
典型用法：
共享内存最常见的用途，就是做数据缓存。比如矩阵乘法中，把子矩阵先加载到共享内存里，然后反复使用，避免反复访问全局内存。
__global__ void matmul_shared(float *A, float *B, float *C, int N) {
__shared__ float As[TILE_SIZE][TILE_SIZE];
__shared__ float Bs[TILE_SIZE][TILE_SIZE];
int bx = blockIdx.x, by = blockIdx.y;
int tx = threadIdx.x, ty = threadIdx.y;
float sum = 0.0f;
for (int t = 0; t < N / TILE_SIZE; t++) {
// 加载到共享内存
As[ty][tx] = A[by * TILE_SIZE + ty][t * TILE_SIZE + tx];
Bs[ty][tx] = B[t * TILE_SIZE + ty][bx * TILE_SIZE + tx];
__syncthreads();
// 计算
for (int k = 0; k < TILE_SIZE; k++)
sum += As[ty][k] * Bs[k][tx];
__syncthreads();
}
C[by * TILE_SIZE + ty][bx * TILE_SIZE + tx] = sum;
}注意：共享内存使用不当会导致bank conflict。我曾经在一个项目中，因为没处理好bank冲突，性能反而比直接用全局内存还差。后来调整了数据布局，才把性能拉回来。
4.3 寄存器（Register）
寄存器是GPU里最快、最私有的存储。每个线程独享，延迟为0（一个时钟周期内完成）。但数量极其有限——每个SM的寄存器总数是固定的，比如1024个或65536个，具体看架构。
关键点：
- 编译器自动分配，程序员无法直接控制
- 寄存器溢出时，数据会被挤到本地内存（Local Memory）
- 寄存器使用量直接影响occupancy（驻留线程数）
我的建议：写kernel时，尽量少用局部变量。不是说不让你用，而是别滥用。每个线程多用几个寄存器，整个SM能同时跑的线程数就少了。这是个权衡——用寄存器换并行度，还是用并行度换寄存器？
避坑指南：我曾经写过一个kernel，里面用了大量临时数组。结果编译器把它们全放到了本地内存里，性能直接腰斩。后来我把数组拆成标量变量，寄存器使用量降下来了，性能也回来了。
4.4 本地内存（Local Memory）
本地内存听起来像“本地”的，其实它物理上就在全局内存里。只是每个线程有自己的私有空间。为什么叫“本地”？因为它只对当前线程可见。
什么时候会用？
- 寄存器不够用时，编译器自动溢出
- 数组索引不是编译期常量时
- 大型结构体或数组
性能影响：本地内存的延迟和全局内存一样高。所以，尽量避免使用。怎么避免？少用大数组，多用标量变量。
小技巧：编译时加上 --ptxas-options=-v 参数，可以查看每个kernel的寄存器使用量和本地内存使用量。我每次调优都会先看这个输出。
4.5 常量内存（Constant Memory）
常量内存是GPU里的一块只读存储，容量很小（通常64KB），但有两个好处：
- 有专门的缓存（Constant Cache）
- 支持广播——同一个warp里的线程访问同一地址时，只需一次读取
适用场景：存储那些不会改变、且所有线程都会用到的数据。比如物理常数、查找表、滤波器系数等。
__constant__ float filter[256];
__global__ void apply_filter(float *input, float *output, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N) {
output[idx] = input[idx] * filter[threadIdx.x % 256];
}
}注意：常量内存的读取速度，取决于warp内所有线程是否访问同一地址。如果每个线程访问不同地址，性能会急剧下降。我见过有人把变量数组放到常量内存里，结果比全局内存还慢。
4.6 纹理内存（Texture Memory）
纹理内存最初是为图形渲染设计的，但后来也被广泛用于通用计算。它有几个独特的能力：
- 硬件插值（线性插值、双线性插值等）
- 边界处理（clamp、wrap、mirror等模式）
- 空间局部性优化（2D/3D访问模式）
什么时候用？
- 图像处理、信号处理
- 需要插值的场景
- 非对齐的随机访问模式
我的经验：有一次做图像缩放，用全局内存做双线性插值，代码写起来很麻烦，性能也一般。后来改用纹理内存，硬件自动做插值，代码简洁了，性能还提升了20%。
注意：纹理内存是只读的。而且绑定纹理需要额外的API调用，代码会稍微复杂一些。不过对于某些场景，这点代价完全值得。
4.7 内存分配与释放
GPU内存的分配和释放，主要通过CUDA API完成。常用的函数有：
| 函数 | 用途 | 说明 |
|---|---|---|
cudaMalloc |
分配全局内存 | 最常用，返回设备指针 |
cudaFree |
释放全局内存 | 必须配对使用 |
cudaMemcpy |
主机与设备间数据拷贝 | 支持同步和异步版本 |
cudaMallocHost |
分配页锁定内存 | 提高主机-设备传输速度 |
cudaMallocManaged |
统一内存分配 | 自动迁移数据，简化编程 |
最佳实践：
- 尽量在程序开始时一次性分配好所有内存，避免频繁分配释放
- 使用
cudaMallocHost分配主机端内存，可以加速PCIe传输 - 统一内存（Managed Memory）虽然方便，但性能不如手动管理
我的习惯：我一般会在代码里加一个内存池。把常用的缓冲区预先分配好，用的时候从池里取，用完归还。这样既避免了频繁的cudaMalloc/cudaFree开销，也减少了内存碎片。
总结一下：GPU内存管理，说白了就是一场“速度与容量”的博弈。全局内存大但慢，共享内存快但小，寄存器最快但最少。你要做的，就是根据数据的使用模式，把数据放到最合适的位置。这个能力，是区分GPU编程新手和老手的关键。
嗯，今天就先聊到这里。下一章我们会讲GPU的同步机制——那又是一个容易踩坑的地方。