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%。

我曾经踩过的坑:全局内存的访问模式必须对齐。如果你让线程随机访问地址,带宽会暴跌。记住:合并访问(coalesced access)是全局内存的命根子。

L2缓存

它是全局内存和L1之间的缓冲。所有SM共享。说实话,你不需要直接控制它,但你要知道它的存在。如果一个数据被多个SM反复读取,L2能帮你省不少带宽。

L1缓存 / 共享内存

这一层最灵活。你可以把它当缓存用(自动管理),也可以当共享内存用(手动管理)。我个人习惯:能用共享内存就别用L1缓存。为什么?因为共享内存的延迟是确定的,而缓存有miss风险。

小技巧:在核函数里用 __shared__ 声明共享内存。记得加 __syncthreads() 同步,不然数据还没写完就被别的线程读了,那结果就乱套了。

常量内存(Constant Memory)

这个层级比较特殊。它只有64KB,但速度极快——前提是同一个warp里的所有线程访问同一个地址。我做过一个光线追踪的项目,把材质参数放在常量内存里,性能提升了30%。

核心要点:常量内存适合存储「所有线程都读同一个值」的数据,比如系数、查找表。如果每个线程读不同的地址,那还不如用全局内存。

局部内存(Local Memory)

这其实是个「陷阱」。当你声明一个局部变量,但寄存器放不下时,编译器会自动把它 spill 到局部内存。说白了,它就是个慢速的备胎。我曾经调试一个核函数,发现性能奇差,一看PTX代码,好家伙,寄存器溢出严重,大量数据被挤到局部内存里了。

避坑指南:别声明太多局部变量。每个线程的寄存器数量有限(通常255个),超了就会 spill。用 --maxrregcount 编译选项可以控制,但别压得太狠,否则 spill 更严重。

寄存器(Registers)

这是金字塔的塔尖。速度最快,没有延迟。每个线程私有。你写的每个局部变量,只要编译器觉得能放,就会放在寄存器里。我建议:尽量多用局部变量,少用数组。因为数组通常会被放到局部内存或共享内存里。

4.4 一张图总结(文字版)

你想象一个金字塔:

  • 塔尖:寄存器(最快,最小,线程私有)
  • 第二层:局部内存(伪快,实际在全局内存)
  • 第三层:常量内存(64KB,只读广播)
  • 第四层:L1/共享内存(手动控制,线程块内共享)
  • 第五层:L2缓存(自动管理,所有SM共享)
  • 塔底:全局内存(最大,最慢,所有线程可见)

记住这个结构,你写CUDA的时候就会下意识地想:我的数据应该放在哪一层? 而不是一股脑全塞全局内存。

5.5 我的个人经验总结

做了这么多年GPU优化,我总结出三条铁律:

  1. 能放寄存器就别放共享内存——寄存器零延迟,共享内存还有bank conflict呢。
  2. 能放共享内存就别放全局内存——共享内存带宽是全局内存的几十倍。
  3. 如果非要用全局内存,一定要合并访问——不然你就是在浪费GPU的带宽。

嗯,说白了,GPU编程就是「数据搬运」的艺术。你搬得越少、搬得越近,性能就越好。下一章咱们会深入讲全局内存的合并访问模式,那才是真正的硬核内容。

课后思考:你的核函数里,哪些数据可以挪到共享内存?哪些变量可能被 spill 到局部内存?打开NVIDIA Nsight看看寄存器使用量,你可能会吓一跳。