3. GPU内存模型详解

各位同学,今天我们来聊聊GPU的内存模型。说实话,我刚开始接触CUDA时,最头疼的就是这堆内存类型。全局内存、共享内存、寄存器、本地内存、常量内存、纹理内存……光名字就够绕的。但你别急,搞懂了它们,你的核函数性能就能起飞。

我个人的习惯是,把GPU内存想象成一个「存储金字塔」。越靠近计算单元的内存,速度越快,但容量越小。反过来,离得远的,容量大,但慢得让人抓狂。咱们一层层来看。

寄存器(Register)

寄存器是离计算单元最近的内存。说白了,它就是每个线程私有的「小本本」。你定义的局部变量,只要编译器觉得够小,就会放在寄存器里。

延迟:几乎为零,1个时钟周期左右。
带宽:极高,每个线程独享。
容量:每个SM(流多处理器)的寄存器总量有限,比如A100每个SM有65536个寄存器。每个线程能分到多少?取决于你开了多少线程。

注意: 我曾经在项目里吃过亏——线程里定义了一大堆局部变量,结果寄存器溢出(register spilling),数据被挤到本地内存里去了,性能直接腰斩。所以,别贪多,一个线程别超过32个寄存器变量,否则编译器会帮你「降级」。

本地内存(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时,总喜欢频繁读写全局内存,结果性能惨不忍睹。后来学乖了——能放共享内存就放共享内存,能合并访问就合并访问。

避坑指南: 我曾经在项目里遇到一个bug——多个线程同时写全局内存的同一个地址,结果数据全乱了。后来加了原子操作(atomicAdd)才解决。但原子操作会串行化,性能会下降。所以,尽量设计成每个线程写不同的地址。

常量内存(Constant Memory)

常量内存是只读的,所有线程都能访问。它也有缓存,但容量很小(64KB)。适合存放不会变的数据,比如权重、系数。

延迟:如果命中缓存,和寄存器差不多;否则和全局内存一样。
带宽:取决于是否命中缓存。
容量:64KB。

我个人习惯用 __constant__ 声明常量内存。注意,常量内存的读取有广播机制——如果同一个warp内的线程读同一个地址,只需要一次内存访问。这个特性在有些场景下很管用。

纹理内存(Texture Memory)

纹理内存最初是为图形学设计的,但现在也用于通用计算。它也是只读的,有专门的缓存和硬件插值单元。

延迟:和常量内存类似,取决于缓存命中。
带宽:对空间局部性好的访问模式,效率很高。
容量:和全局内存一样大。

纹理内存有个好处——它支持「边界处理」和「硬件插值」。比如做图像处理时,你访问的坐标超出边界,它会自动返回边界值或重复值。我在做卷积神经网络时用过纹理内存,效果不错。

延迟与带宽对比表

内存类型 延迟(时钟周期) 带宽(相对值) 容量 作用域
寄存器 ~1 极高 每个SM有限 线程私有
本地内存 400-800 低(但有缓存) 每个线程512KB 线程私有
共享内存 20-30 高 每个SM 164KB 线程块内共享
全局内存 400-800 高(但延迟大) 几十GB 所有线程
常量内存 ~1(缓存命中) 中等 64KB 所有线程(只读)
纹理内存 ~1(缓存命中) 中等 同全局内存 所有线程(只读)

知识体系结构图

GPU内存模型金字塔 全局内存 (Global Memory) 容量大 · 延迟高 · 所有线程可读写 常量/纹理内存 (Constant/Texture) 只读 · 有缓存 · 广播/插值特性 共享内存 (Shared Memory) 片上 · 低延迟 · 线程块内共享 本地内存 (Local) 寄存器溢出 · 线程私有 寄存器 最快 · 线程私有 速度 ↑ 容量 ↑

这张图很直观地展示了内存的层级关系。越往上,速度越快,容量越小。写核函数时,你要做的就是尽量把数据往上层搬。

总结一下

GPU内存模型的核心就一句话:让数据尽可能靠近计算单元。寄存器最快,但容量小;共享内存次之,但需要手动管理;全局内存最慢,但容量大。常量内存和纹理内存是「特化」的只读内存,适合特定场景。

我个人建议,写核函数时先想清楚数据流:哪些数据是每个线程私有的?哪些是线程块内共享的?哪些是全局只读的?想清楚了,内存类型自然就选对了。