1、GPU内存层次结构全景:从全局内存到寄存器,一张图看懂GPU存储金字塔
各位同学,咱们今天聊聊GPU的内存。说实话,我见过太多人写CUDA代码,一上来就猛怼全局内存,结果性能惨不忍睹。其实啊,GPU的内存不是铁板一块,它是个金字塔结构。你想想看,从最慢的全局内存到最快的寄存器,速度差了上百倍。不理解这个层次结构,你写的核函数大概率跑不满。
1.1 为什么需要内存层次结构?
先问个问题:为什么GPU不直接用一块又大又快的内存?答案很简单——成本和物理限制。我当年在优化一个流体力学模拟时,就吃过这个亏。当时我天真地以为把所有数据塞进全局内存就行,结果带宽成了瓶颈,GPU利用率不到20%。
GPU的内存设计,说白了就是个「妥协的艺术」:
- 容量要大——装得下海量数据
- 速度要快——喂饱几千个计算核心
- 成本要低——不然一块卡卖你几十万
这三者不可能同时满足。所以NVIDIA搞了个金字塔:越往上越快、越贵、容量越小;越往下越慢、越便宜、容量越大。
2.2 一张图看懂GPU存储金字塔
我个人习惯把GPU内存分成六个层级,从下往上依次是:
| 层级 | 名称 | 典型容量 | 相对速度 | 作用域 |
|---|---|---|---|---|
| 6(最慢) | 全局内存 | 8-80 GB | 1x(基准) | 所有线程 |
| 5 | L2缓存 | 几MB | ~10x | 所有线程 |
| 4 | L1缓存/共享内存 | 48-128 KB | ~50x | 线程块内 |
| 3 | 常量内存 | 64 KB | ~100x(只读广播) | 所有线程 |
| 2 | 局部内存 | 线程私有 | ~1x(实际在全局内存) | 单个线程 |
| 1(最快) | 寄存器 | 每个SM几千个 | ~500x | 单个线程 |
嗯,这里要注意:局部内存虽然名字带「局部」,但它实际上是在全局内存里划了一块。我刚开始学的时候就被这个名字骗过,以为它很快,结果一测延迟高得吓人。
3.3 每个层级的特点与避坑指南
全局内存(Global Memory)
这是GPU的「大仓库」。所有线程都能读写,容量最大,但速度最慢。我在项目中遇到过一个问题:一个矩阵乘法核函数,直接读写全局内存,带宽利用率只有30%。后来改成共享内存分块,直接飙到85%。
L2缓存
它是全局内存和L1之间的缓冲。所有SM共享。说实话,你不需要直接控制它,但你要知道它的存在。如果一个数据被多个SM反复读取,L2能帮你省不少带宽。
L1缓存 / 共享内存
这一层最灵活。你可以把它当缓存用(自动管理),也可以当共享内存用(手动管理)。我个人习惯:能用共享内存就别用L1缓存。为什么?因为共享内存的延迟是确定的,而缓存有miss风险。
__shared__ 声明共享内存。记得加 __syncthreads() 同步,不然数据还没写完就被别的线程读了,那结果就乱套了。
常量内存(Constant Memory)
这个层级比较特殊。它只有64KB,但速度极快——前提是同一个warp里的所有线程访问同一个地址。我做过一个光线追踪的项目,把材质参数放在常量内存里,性能提升了30%。
局部内存(Local Memory)
这其实是个「陷阱」。当你声明一个局部变量,但寄存器放不下时,编译器会自动把它 spill 到局部内存。说白了,它就是个慢速的备胎。我曾经调试一个核函数,发现性能奇差,一看PTX代码,好家伙,寄存器溢出严重,大量数据被挤到局部内存里了。
--maxrregcount 编译选项可以控制,但别压得太狠,否则 spill 更严重。
寄存器(Registers)
这是金字塔的塔尖。速度最快,没有延迟。每个线程私有。你写的每个局部变量,只要编译器觉得能放,就会放在寄存器里。我建议:尽量多用局部变量,少用数组。因为数组通常会被放到局部内存或共享内存里。
4.4 一张图总结(文字版)
你想象一个金字塔:
- 塔尖:寄存器(最快,最小,线程私有)
- 第二层:局部内存(伪快,实际在全局内存)
- 第三层:常量内存(64KB,只读广播)
- 第四层:L1/共享内存(手动控制,线程块内共享)
- 第五层:L2缓存(自动管理,所有SM共享)
- 塔底:全局内存(最大,最慢,所有线程可见)
记住这个结构,你写CUDA的时候就会下意识地想:我的数据应该放在哪一层? 而不是一股脑全塞全局内存。
5.5 我的个人经验总结
做了这么多年GPU优化,我总结出三条铁律:
- 能放寄存器就别放共享内存——寄存器零延迟,共享内存还有bank conflict呢。
- 能放共享内存就别放全局内存——共享内存带宽是全局内存的几十倍。
- 如果非要用全局内存,一定要合并访问——不然你就是在浪费GPU的带宽。
嗯,说白了,GPU编程就是「数据搬运」的艺术。你搬得越少、搬得越近,性能就越好。下一章咱们会深入讲全局内存的合并访问模式,那才是真正的硬核内容。