GPU内存模型详解:全局内存、共享内存、寄存器、本地内存、常量内存、纹理内存
各位同学,今天我们来聊聊GPU的内存模型。说实话,我刚开始接触CUDA时,最头疼的就是这堆内存类型。全局内存、共享内存、寄存器、本地内存、常量内存、纹理内存……光名字就够绕的。但你别急,搞懂了它们,你的核函数性能就能起飞。
我个人的习惯是,把GPU内存想象成一个「存储金字塔」。越靠近计算单元的内存,速度越快,但容量越小。反过来,离得远的,容量大,但慢得让人抓狂。咱们一层层来看。
寄存器(Register)
寄存器是离计算单元最近的内存。说白了,它就是每个线程私有的「小本本」。你定义的局部变量,只要编译器觉得够小,就会放在寄存器里。
延迟:几乎为零,1个时钟周期左右。
带宽:极高,每个线程独享。
容量:每个SM(流多处理器)的寄存器总量有限,比如A100每个SM有65536个寄存器。每个线程能分到多少?取决于你开了多少线程。
本地内存(Local Memory)
本地内存这个名字容易误导人。它其实不是「本地」的,而是存在全局内存里的。只是每个线程觉得自己有块私有空间。
什么时候会用上?寄存器不够用的时候。比如你定义了一个大数组,或者编译器觉得变量太多,就会把一部分数据放到本地内存。
延迟:和全局内存一样,几百个时钟周期。
带宽:和全局内存一样,但因为有缓存,稍微好点。
容量:每个线程最多512KB(取决于GPU架构)。
嗯,这里要注意:本地内存是性能杀手。我建议你写代码时,用 --ptxas-options=-v 编译选项看看寄存器使用情况。如果看到「local memory」字样,赶紧优化。
共享内存(Shared Memory)
共享内存是我最喜欢的内存类型。它位于芯片上,同一个线程块(block)内的所有线程都能访问。速度比全局内存快得多。
延迟:大约20-30个时钟周期。
带宽:极高,是全局内存的10倍以上。
容量:每个SM有限,比如A100是164KB(可配置)。
你想想看,为什么共享内存这么快?因为它就在SM内部,没有走片外总线。我在做矩阵乘法优化时,就是靠共享内存把数据「搬」到片上,然后反复利用,性能提升了5倍。
__shared__ 声明,然后用 __syncthreads() 同步。记住,不同步的话,数据竞争会让你怀疑人生。
全局内存(Global Memory)
全局内存是GPU的「主存」。所有线程都能读写,但速度最慢。它位于显存里,通过PCIe或NVLink与CPU通信。
延迟:400-800个时钟周期。
带宽:很高,比如A100有2TB/s,但延迟太大。
容量:最大,通常几十GB。
为什么全局内存这么慢?因为它离计算单元太远了。数据要从显存芯片传到SM,中间经过很多层级。我刚开始写CUDA时,总喜欢频繁读写全局内存,结果性能惨不忍睹。后来学乖了——能放共享内存就放共享内存,能合并访问就合并访问。
常量内存(Constant Memory)
常量内存是只读的,所有线程都能访问。它也有缓存,但容量很小(64KB)。适合存放不会变的数据,比如权重、系数。
延迟:如果命中缓存,和寄存器差不多;否则和全局内存一样。
带宽:取决于是否命中缓存。
容量:64KB。
我个人习惯用 __constant__ 声明常量内存。注意,常量内存的读取有广播机制——如果同一个warp内的线程读同一个地址,只需要一次内存访问。这个特性在有些场景下很管用。
纹理内存(Texture Memory)
纹理内存最初是为图形学设计的,但现在也用于通用计算。它也是只读的,有专门的缓存和硬件插值单元。
延迟:和常量内存类似,取决于缓存命中。
带宽:对空间局部性好的访问模式,效率很高。
容量:和全局内存一样大。
纹理内存有个好处——它支持「边界处理」和「硬件插值」。比如做图像处理时,你访问的坐标超出边界,它会自动返回边界值或重复值。我在做卷积神经网络时用过纹理内存,效果不错。
延迟与带宽对比表
| 内存类型 | 延迟(时钟周期) | 带宽(相对值) | 容量 | 作用域 |
|---|---|---|---|---|
| 寄存器 | ~1 | 极高 | 每个SM有限 | 线程私有 |
| 本地内存 | 400-800 | 低(但有缓存) | 每个线程512KB | 线程私有 |
| 共享内存 | 20-30 | 高 | 每个SM 164KB | 线程块内共享 |
| 全局内存 | 400-800 | 高(但延迟大) | 几十GB | 所有线程 |
| 常量内存 | ~1(缓存命中) | 中等 | 64KB | 所有线程(只读) |
| 纹理内存 | ~1(缓存命中) | 中等 | 同全局内存 | 所有线程(只读) |
知识体系结构图
这张图很直观地展示了内存的层级关系。越往上,速度越快,容量越小。写核函数时,你要做的就是尽量把数据往上层搬。
总结一下
GPU内存模型的核心就一句话:让数据尽可能靠近计算单元。寄存器最快,但容量小;共享内存次之,但需要手动管理;全局内存最慢,但容量大。常量内存和纹理内存是「特化」的只读内存,适合特定场景。
我个人建议,写核函数时先想清楚数据流:哪些数据是每个线程私有的?哪些是线程块内共享的?哪些是全局只读的?想清楚了,内存类型自然就选对了。
好了,这一章就到这里。记住,内存优化是GPU性能调优的基石。下一章我们会深入共享内存的bank conflict问题,那才是真正考验功底的地方。