3、内存子系统剖析：全局内存、共享内存、寄存器文件与缓存层级
内存子系统，说白了就是GPU的「粮草官」。程序跑得快不快，很多时候不是计算单元不够猛，而是数据没喂到嘴边。我这些年调过的性能瓶颈，十有八九都跟内存沾边。今天咱们就把全局内存、共享内存、寄存器文件和缓存层级这四块，掰开了揉碎了讲清楚。
3.1 全局内存：最宽敞但也最慢的「大仓库」
全局内存是GPU里容量最大的存储空间，所有线程都能访问。但容量大意味着距离远，延迟高。我记得刚入行时，有个同事写了个矩阵乘法，数据全放全局内存，结果跑起来比CPU还慢。他一脸懵地问我：「GPU不是号称并行计算吗？」
嗯，问题就出在全局内存的访问模式上。
核心要点：全局内存的带宽利用率，取决于合并访问（Coalesced Access）。
什么叫合并访问？简单说，就是让同一个warp里的32个线程，访问连续的内存地址。GPU硬件会把这32个请求合并成一次或几次大的内存事务。你想想看，一次取128字节，跟32次各取4字节，效率能一样吗？
我给大家看个对比：
// ❌ 非合并访问：线程0访问A[0][0]，线程1访问A[1][0]...
// 地址不连续，性能极差
float val = A[threadIdx.y][threadIdx.x];
// ✅ 合并访问：线程0访问A[0][0]，线程1访问A[0][1]...
// 地址连续，一次事务搞定
float val = A[threadIdx.x][threadIdx.y];我在项目中遇到过最典型的案例，是一个图像滤波算子。原始代码按列处理，全局内存访问全是跨步的。改成按行处理后，带宽利用率从15%直接飙到85%，整体性能提升了4倍多。
个人习惯：写全局内存访问代码时，我总会先画个线程-地址映射图。花5分钟画图，省5小时调优。
3.2 共享内存：片上的「高速缓存池」
共享内存是GPU架构里最性感的特性之一。它位于芯片内部，延迟比全局内存低一个数量级。但容量有限——通常每个SM只有几十到一百多KB。
共享内存的真正价值在于：线程间数据共享和数据重用。
举个例子，做3x3的卷积操作。每个输出像素需要周围9个输入像素。如果直接从全局内存读，每个像素会被相邻的多个线程重复读取。这时候共享内存就派上用场了：
__global__ void conv3x3(float* in, float* out, int W, int H) {
__shared__ float tile[TILE_SIZE + 2][TILE_SIZE + 2]; // 带halo区域
int tx = threadIdx.x, ty = threadIdx.y;
int gx = blockIdx.x * TILE_SIZE + tx - 1;
int gy = blockIdx.y * TILE_SIZE + ty - 1;
// 协作加载数据到共享内存
if (gx >= 0 && gx < W && gy >= 0 && gy < H)
tile[ty][tx] = in[gy * W + gx];
else
tile[ty][tx] = 0.0f;
__syncthreads(); // 同步！确保所有数据加载完成
// 从共享内存计算
if (tx >= 1 && tx <= TILE_SIZE && ty >= 1 && ty <= TILE_SIZE) {
float sum = 0.0f;
for (int dy = -1; dy <= 1; dy++)
for (int dx = -1; dx <= 1; dx++)
sum += tile[ty + dy][tx + dx] * kernel[dy+1][dx+1];
out[(gy) * W + (gx)] = sum;
}
}这里有个坑，我曾经踩过——bank conflict。共享内存被分成32个bank，每个bank宽度4字节。如果同一个warp里的多个线程访问同一个bank的不同地址，就会发生冲突，导致访问串行化。
避坑指南：我曾经花了两天时间排查一个性能问题，最后发现是共享内存的步长正好是4字节的整数倍，导致所有线程都扎堆在bank 0上。解决办法很简单——在声明时加一列padding：__shared__ float tile[32][32 + 1];
3.3 寄存器文件：最快但最稀缺的「贴身口袋」
寄存器是GPU里最快的存储，没有之一。每个SM有大量的寄存器（比如A100有65536个），但分摊到每个线程头上，其实很有限。
为什么说寄存器是稀缺资源？因为寄存器用量直接决定了SM上能同时驻留的线程数量（occupancy）。
| 每个线程寄存器数 | 每个SM最大线程数 | Occupancy |
|---|---|---|
| 32 | 2048 | 100% |
| 64 | 1024 | 50% |
| 128 | 512 | 25% |
你看，寄存器用多了，occupancy就往下掉。但occupancy低就一定慢吗？不一定。我见过一些计算密集型的kernel，即使occupancy只有25%，因为每个线程有充足的寄存器来缓存中间结果，性能反而更好。
这里有个权衡：寄存器溢出（register spilling）。当编译器发现寄存器不够用时，会把变量「挤」到局部内存里——而局部内存其实就在全局内存里，延迟瞬间爆炸。
我的调试习惯：编译时加--ptxas-options=-v参数，看编译器报告里有没有「spill」字样。如果有，说明寄存器不够用了。这时候要么减少每个线程的变量数，要么用__launch_bounds__限制最大线程数来给编译器更多寄存器空间。
3.4 缓存层级：看不见的「性能加速器」
现代GPU的缓存体系越来越复杂。以NVIDIA的架构为例，有L1缓存、L2缓存，还有只读缓存（纹理缓存）。
L1缓存和共享内存其实是共用同一块物理SRAM。你可以通过API调整它们的比例：
// 分配48KB共享内存 + 16KB L1缓存
cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, 48 * 1024);
// 或者反过来，16KB共享内存 + 48KB L1缓存
// 默认配置通常就够了，除非你有特殊需求L2缓存是所有SM共享的。我记得有一次调优，发现某个kernel的L2缓存命中率只有30%。排查后发现是数据访问模式太随机，导致缓存行频繁被踢出。后来改成了分块处理，让每个数据块在L2里「住」得久一点，命中率提到了70%以上。
缓存命中的重要性，我给大家算笔账：
- L1命中：~30个周期
- L2命中：~200个周期
- 全局内存：~400-800个周期
你想想看，如果能把一次全局内存访问变成L1命中，速度能快20倍。这就是为什么我们要花那么多心思在数据局部性上。
3.5 实战总结：内存子系统的选择策略
说了这么多，到底什么时候用哪种内存？我给大家一个简单的决策树：
- 数据只被一个线程用一次 → 全局内存（合并访问）
- 数据被多个线程反复用 → 共享内存（注意bank conflict）
- 每个线程的临时变量 → 寄存器（别溢出）
- 只读数据且访问模式规则 → 纹理缓存/只读缓存
最后说一句：别迷信理论带宽。我见过太多人拿着800GB/s的带宽参数，实际跑出来只有100GB/s。原因无他——访问模式不对。内存调优的本质，就是让硬件的每个存储层级都物尽其用。这个道理，你越早明白越好。