想象一座繁忙的“芯片城市”:GPU 里有成千上万个计算小分队——线程,它们像流水线上的工人一样不停地做算术。做算术前要有数据,这就像餐厅做菜前要有食材。食材放在哪里、离灶台有多远,直接决定了出菜速度。GPU 的显存层次就是一套分级“仓储物流系统”:工人腰间的工具袋(寄存器)、同组共用的工作台(共享内存)、街角的便利店和社区超市(L1/L2 缓存),以及郊区的巨型中央仓库(全局内存)。腰包最快也最小,中央仓库最大但来回一趟要好几百个周期。这篇文章带你沿这套物流系统走一圈:每一级有什么特点、为什么延迟差了好几个数量级,以及如何像资深“仓管员”一样优化搬运路线,让深度学习算子跑得飞快。

岗位关联

对 AI 从业者而言,显存层次是读懂 GPU 性能瓶颈的“透视镜”。无论是训练大模型、写 CUDA 算子,还是做推理加速,90% 的性能问题都藏在“数据放错层级”里。

目录
  1. 可视化演示
  2. 层次总览:芯片城市的四级仓储
  3. 寄存器:工人腰间的工具袋
  4. 共享内存:车间公用的工作台
  5. 缓存层次:街角便利店与社区超市
  6. 全局内存:郊区的中央仓库
  7. 局部性原理:熟悉的人和邻居
  8. 优化策略:仓管员的搬运艺术
  9. 深度学习实践:大模型时代的显存兵法
  10. AI 应用场景
  11. Python 演示
  12. 常见误区
  13. 面试题

可视化演示

展示NPU(神经网络处理器)核心的脉动阵列架构,包含MAC计算单元数据流、输入权重加载过程,以及NPU与GPU/CPU在AI推理场景下的性能对比。

本节配有交互动效(进入视口后自动加载,约 3KB)。

NPU神经网络处理器架构 脉动阵列:MAC单元矩阵高效执行矩阵乘法 数据流:输入与权重同步流入PE阵列 INT8量化:低精度高吞吐推理加速 NPU = Systolic Array + INT8 MAC + 专用存储层次

层次总览:芯片城市的四级仓储

想象一座芯片城市,计算单元(SM)是繁忙的车间,每个车间里有几十个“工人线程”。他们干活需要原料——数据。原料放在哪里,决定了工人能否流畅干活。于是城市建了四层仓储:工人腰间工具袋、车间公用工作台、街角便利店社区超市、以及郊区中央仓库。这四级仓储从近到远、从贵到便宜、从快到慢,彼此配合构成了显存层次。

GPU 的显存层次可以粗分两类:片上存储(寄存器、共享内存、L1/L2 缓存)和片外存储(全局内存)。凡是物理上离计算单元越近的层级,容量越小、造价越贵、速度越快;离得越远,容量越大、速度越慢。程序员最关心的是“一次访存要等多少个周期”,而这串数字跨越了三个数量级。

所以 GPU 程序优化的本质,就是尽可能让工人使用腰包和工作台,少去中央仓库。理解了这张“价格表”,你就握住了显存优化的钥匙。

Treg ≪ Tshared≈TL1< TL2 ≪ Tglobal\displaystyle T_{reg} \ \ll \ T_{shared} \approx T_{L1} < \ T_{L2} \ \ll \ T_{global}
这条不等式给出了 GPU 访存延迟的“价格表”:寄存器几乎免费,全局内存贵得惊人。
  1. 左侧最快$T_{reg}$ 是寄存器延迟,约 1 个周期,相当于工人伸手从腰包取工具。
  2. 中间过渡带$T_{shared} \approx T_{L1}$ 表示共享内存和 L1 缓存延迟同处一个量级,约 20~35 周期,像去几步外的公用工作台或街角便利店。
  3. L2 缓存更远$T_{L2}$ 约为 200 周期,像开车去社区超市,全芯片所有车间共用。
  4. 右侧最慢$T_{global}$ 约为 400~800 周期,像从中央仓库调货,来回一趟耗时惊人。

寄存器:工人腰间的工具袋

每个工人线程腰间挂着一个迷你工具袋——寄存器文件。它是最私密、最快速的存储:只有自己能用,别人碰不到。GPU 为每个线程分配若干寄存器,访问它们只要 1 个周期,几乎不影响性能。但工具袋的容量极其有限:一个线程通常只能分到几十个寄存器,超过限制就被“挤出”或降速。

寄存器文件在每个 SM 里总容量固定,比如 65536 个 32 位寄存器。每个线程用多少寄存器,直接决定一个 SM 能同时驻留多少线程:

举例:65536 个寄存器,每人 32 个,能驻留 2048 线程;如果 kernel 写得寄存器爆炸、每人用 128 个,驻留线程骤降到 512,延迟就无法被大量线程掩盖。所以优化要“抠门”:少用寄存器,把资源还给并行度。

Nthread=Nreg_fileNreg_thread\displaystyle N_{thread} = \frac{N_{reg\_file}}{N_{reg\_thread}}
这是“寄存器预算”公式:每人少拿一点,就能让更多工人同时上岗。
  1. 分子是总预算$N_{reg\_file}$ 表示每个 SM 的寄存器总数,它是固定的硬件资源。
  2. 分母是人均开销$N_{reg\_thread}$ 表示每个线程实际占用的寄存器数量,由编译器决定。
  3. 整体是驻留人数$N_{thread}$ 是 SM 能同时容纳的线程数,也叫影响“占用率”的关键指标。

共享内存:车间公用的工作台

同一车间的工人们(一个线程块)身边有一张大家共用的工作台——共享内存。它由程序员显式管理,擅长让线程之间交换数据、协作完成同一项任务。共享内存访问延迟约 20 个周期,远快于全局内存,但每 SM 只有几十 KB,容量很紧张。

这张工作台的奇特之处在于它有 32 个“抽屉”(存储体,bank)。如果 32 个线程同时访问 32 个不同抽屉的数据,一次就能完成;但如果两个线程恰好访问同一抽屉的两个不同地址,就会发生存储体冲突,访问被迫串行执行:

所以用共享内存时,要精心设计索引,让同一时刻访问的地址分散在 32 个银行中,避免挤在同一个抽屉里排队。

效率=1nconflict\displaystyle \text{效率} = \frac{1}{n_{conflict}}
冲突率直接打折:冲突越多,工作台的并行优势越小。
  1. 分子是一次理想访问理想情况下 $n_{conflict}=1$,效率为 100%,32 个数据一趟取完。
  2. 分母是冲突线程数$n_{conflict}$ 表示落在同一个银行的不同地址个数,例如 2 表示有两个线程挤同一抽屉。
  3. 公式含义2 路冲突时效率降到 $\frac{1}{2}$,一次访问变成两趟,速度减半。

缓存层次:街角便利店与社区超市

芯片城市里还开着两家“自动补货”的超市:L1 缓存是每个车间门口的街角便利店,容量几十 KB,延迟约 30 周期;L2 缓存是服务全城的社区超市,容量几 MB,延迟约 200 周期。它们都由硬件自动管理,不需要程序员显式搬运,配合“缓存行”机制把连续数据成批拉近。

缓存的价值在于命中与未命中。命中了,数据在近处直接取用;未命中,就要去更远的层级搬一次。整体平均访存延迟可以这样估算:

理解缓存的“自动补货”逻辑后,你就明白为什么要写“对缓存友好”的代码:多访问同一块数据、多访问连续地址,让便利店多命中几次,少跑社区超市。

Tavg=h⋅Thit+(1−h)⋅Tmiss\displaystyle T_{avg} = h \cdot T_{hit} + (1-h) \cdot T_{miss}
命中率越高,平均延迟越接近“近处”的速度。
  1. 命中率 h$h$ 表示访问在缓存中直接命中的概率,例如 $h = 90\%$。
  2. 命中与未命中加权$T_{hit}$ 是命中延迟(如 30 周期),$T_{miss}$ 是未命中去下级取回的延迟(如 200 周期)。
  3. 代入示例$T_{avg} = 0.9 \times 30 + 0.1 \times 200 = 47$ 周期,比每次都去全局内存快得多。

全局内存:郊区的中央仓库

全局内存在 GPU 的显存层次中扮演“终极大仓库”的角色:容量最大(几十 GB),价格最便宜,但延迟也最高(400~800 周期)。GPU 与全局内存之间有高带宽总线,比如 HBM 可以有每秒数 TB 的带宽,但前提是你能“装满每一辆卡车”——这就是合并访问。

现代 GPU 按 32 字节扇区从全局内存取数据。如果一个线程束(32 个线程)访问的地址是连续的、对齐的,硬件就能用最少的事务把数据一次搬回。反之,如果线程东跳一个、西跳一个,同一辆卡车只装了几件货,效率自然暴跌:

写核函数时,尽量让相邻线程访问相邻地址,让每次全局内存访问都像“满载卡车”一样高效。这是 GPU 编程中最经典的性能口诀之一。

η=BusedBfetched=n×4⌈n×432⌉×32\displaystyle \eta = \frac{B_{used}}{B_{fetched}} = \frac{n \times 4}{\left\lceil \frac{n \times 4}{32} \right\rceil \times 32}
合并访问的效率 η 越接近 1,表示搬运的每个字节都没浪费。
  1. 使用字节数$n \times 4$ 表示 $n$ 个线程每个读取 4 字节(float),共用了多少有效数据。
  2. 搬运字节数$\lceil \frac{n \times 4}{32} \rceil \times 32$ 表示硬件按 32 字节扇区向上取整,实际搬回多少字节。
  3. 效率计算32 个线程连续访问 128 字节时,$\eta = \frac{128}{128} = 1$,满载而归。

局部性原理:熟悉的人和邻居

缓存和工作台之所以快,靠的是两条朴素的规律:时间局部性——刚用过的东西很可能马上再用;空间局部性——用了一个数据,旁边的数据大概率也要用。就像工人喜欢把常用工具放在手边,买东西会一次把相邻货架扫一遍。

局部性决定了程序属于“计算密集”还是“内存密集”。衡量这一性质的核心指标是算术强度:

掌握算术强度,你就拥有了判断“瓶颈在哪”的第一把尺子:搬数据的时间远大于算数据的时间时,任何计算优化都是杯水车薪。

I=FQ\displaystyle I = \frac{F}{Q}
算术强度 I 表示“每搬 1 字节数据,做了多少次运算”。
  1. 分子是运算量$F$ 是程序完成的总浮点运算次数(FLOPs),例如两个 $N \times N$ 矩阵相乘约为 $2N^3$。
  2. 分母是搬运量$Q$ 是访问全局内存的字节数,例如每个矩阵元素 4 字节。
  3. 强度分类$I$ 很小(如向量加法约 0.25)属于内存密集;$I$ 很大(如大矩阵乘法随 $N$ 增长)属于计算密集。

优化策略:仓管员的搬运艺术

一个经验丰富的仓管员不会让工人每次都跑去中央仓库取一件货,而是会统筹规划:先把一批货搬到车间工作台,让工人就地取用多次。这套思路对应 GPU 优化的两大招:数据复用和分块。把要反复使用的数据切成小块载入共享内存,每次处理一小块,积少成多把全局流量压低几十倍。

那到底该优化计算还是优化访存?用 Roofline 模型一画便知:横轴是算术强度,纵轴是可达性能。性能不会超过天花板——计算峰值和带宽上限的交界处:

分块后数据留在共享内存里反复复用,$I$ 上升,程序从陡坡段爬向平台段。这就是仓管员的智慧:与其跑一百趟,不如一次搬一箱。

P=min⁡(Ppeak, I×BWpeak)\displaystyle P = \min\left(P_{peak},\ I \times BW_{peak}\right)
实际算术强度落在陡坡段就是内存受限,落在平台段就是计算受限。
  1. 内存受限分支$I \times BW_{peak}$ 表示数据搬运速度决定的性能上限,当 $I$ 较小时这个数更低。
  2. 计算受限分支$P_{peak}$ 表示硬件算力峰值,当 $I$ 足够大时空算力成为瓶颈。
  3. 取最小值$\min\left(\cdot, \cdot\right)$ 表示最终性能受两个瓶颈中更矮的那个限制。

深度学习实践:大模型时代的显存兵法

深度学习的主力运算是矩阵乘法和卷积,它们的共同特点是数据复用潜力巨大。一个 $N \times N$ 的矩阵乘法,若不优化要反复读写全局内存;用分块策略后,每个数据块在共享内存中被复用 $N/tile$ 次,全局搬运量大幅下降。其算术强度近似:

这正是 FlashAttention 等高效算子的根基:把大矩阵切成小块,让数据尽量停留在共享内存,把全局内存的访问量从 $O(N^2)$ 降到 $O(N)$。训练大模型时,显存层次的运用已然是兵家必争之地——谁的数据离计算更近,谁就能跑得更快。

AIGEMM=2N33N2=2N3\displaystyle AI_{GEMM} = \frac{2N^3}{3N^2} = \frac{2N}{3}
矩阵越大,算术强度越高,越倾向计算受限;小矩阵则容易掉进带宽受限区。
  1. 分子是计算量$2N^3$ 是矩阵乘法的浮点运算次数,每个输出元素要做 $2N$ 次乘加。
  2. 分母是搬运量$3N^2$ 是读取 $A$、$B$ 并写回 $C$ 的最小字节数。
  3. 化简结果$\frac{2N}{3}$ 表示当 $N=1024$ 时约 683 FLOPS/字节,远高于当前显存带宽所能喂饱的水平,因此大 GEMM 是计算受限。

AI 应用场景

场景一:深度学习算子优化(训练/推理)

阶段:模型训练与推理加速

解决的问题:PyTorch 默认算子对特定形状的张量往往不是最优实现,频繁把大矩阵从全局内存搬进搬出,导致 GPU 算力空转。利用共享内存分块与合并访问,可以让算子更贴合显存层次。

具体案例:某 NLP 模型训练时,Self-Attention 的 QKV 计算被重写为类似 FlashAttention 的分块算子:Q、K、V 以小块形式载入共享内存,全局内存往返次数从 $O(N^2)$ 降到 $O(N)$,端到端训练速度提升约 1.8 倍。

场景二:大模型 KV Cache 推理加速

阶段:大模型推理服务优化

解决的问题:自回归生成时每步都要读取越来越长的 KV Cache,这些数据驻留在大容量全局内存中。若不注意缓存局部性,每一步都会付出高额延迟。

具体案例:推理框架把 KV Cache 按固定 block 管理,并让查询向量以合并访问方式读取连续缓存行;热数据尽量留在 L2 缓存中,使平均每次 token 生成的访存延迟降低 40%。

场景三:AI 芯片设计中的存储规划

阶段:芯片架构设计验证

解决的问题:设计 AI 芯片时,片上 SRAM(共享内存/缓存)面积昂贵,不能无限加大;分配多少给 L1、多少给共享内存,需要结合目标算子的算术强度做权衡。

具体案例:芯片设计团队用 Roofline 模型分析主流视觉模型的分块策略,确定每 SM 配置 192KB 共享内存和 128KB L1 缓存;编译器依据模型自动选择 tile 大小,使矩阵乘法的片上数据复用率提升 3 倍。

Python 演示

python35 行
def matmul_traffic(N, tile):    """估算矩阵乘法的 DRAM 流量(字节)。     参数:        N: 矩阵边长(N x N)        tile: 分块边长;0 表示朴素实现(不分块)    返回:        理论 DRAM 流量(字节)    """    if tile == 0:        # 朴素实现:每个 A/B 元素都被全局读写 O(N) 次        a_flow = N ** 3        b_flow = N ** 3        c_flow = N ** 2    else:        # 分块实现:A/B 各被复用 N/tile 次        a_flow = N ** 3 / tile        b_flow = N ** 3 / tile        c_flow = N ** 2    return (a_flow + b_flow + c_flow) * 4  # float32 每元素 4 字节  def main():    N = 1024    print(f"N={N} 的矩阵乘法,对比不同分块大小:")    for tile in [0, 16, 64]:        bytes_flow = matmul_traffic(N, tile)        gb = bytes_flow / 1e9        # 假设可用带宽为 1 TB/s,计算理想耗时        ms = gb / 1.0 * 1000        print(f"tile={tile:>2}  DRAM流量 {gb:8.2f} GB, 理想耗时 {ms:8.2f} ms")  if __name__ == "__main__":    main()
目的:演示数据分块复用如何把矩阵乘法的全局内存流量从 8GB 级别压到 0.1GB 级别,解释为何 GPU 优化离不开显存层次。
  1. a_flow = N ** 3朴素实现中每个 A 元素在 k 循环里被反复读取 N 次,总量为 N³。
  2. a_flow = N ** 3 / tile分块后 A 的每个子块被载入共享内存一次、复用 tile² 次,流量下降约 tile 倍。
  3. (a_flow + b_flow + c_flow) * 4将“元素访问次数”换算成字节数(float32 占 4 字节)。
  4. ms = gb / 1.0 * 1000假设 1TB/s 可用带宽,把流量换算成理想耗时,直观量化优化收益。

常见误区

误区我以为共享内存既然在芯片上,容量怎么也有几十 MB,把大数组全放进去没问题。结果程序一运行,kernel 直接启动失败。

正解共享内存每个 SM 通常只有几十到一百多 KB,比 L2 缓存小得多。它像一张工作台,桌面面积有限,只适合放“反复使用的小块数据”。大数组应该留在全局内存,通过分块逐步搬上工作台。

为何常见:因为 CPU 的 L2 缓存动辄几 MB,容易让人高估 GPU 片上存储的容量。

误区既然全局内存要几百个周期,那我少读它就行,哪怕让很多线程重复计算也行。

正解全局内存并非“碰不得”,它拥有整个显存系统中最大的容量和极高的带宽。关键在于用合并访问把每次搬运装满,同时配合共享内存减少总访问次数。跳过全局内存反而可能引入重复计算,得不偿失。

为何常见:初学者把“延迟高”误解为“带宽低”,其实全局内存的总带宽高达数 TB/s,问题是延迟和浪费。

误区我把很多中间结果都存在寄存器里,认为这样能减少访存、加速运行。结果性能反而下降。

正解寄存器文件总量固定,线程占用越多,能驻留的线程数就越少,GPU 用大量线程隐藏延迟的优势被削弱。合理做法是控制寄存器数量在编译器默认范围内,必要时用 __launch_bounds__ 提示编译器平衡占用率与寄存器溢出。

为何常见:看到“寄存器最快”后容易片面追求寄存器存储,忽略了资源总量约束。

误区我做深度学习,用 PyTorch 写模型,GPU 显存层次是游戏开发者才关心的事。

正解深度学习核心的矩阵乘法、卷积和注意力算子,性能几乎都由访存模式决定。FlashAttention、算子融合、kernel 调优都建立在显存层次的知识上。不管前端是 PyTorch 还是 TensorFlow,底层 CUDA 算子都在精打细算每一层缓存。

为何常见:框架封装让开发者看不到底层算子,误以为框架层已经“优化完了”。

面试题

Q1. GPU 显存层次中,访问延迟最低的是哪一级?

单选题 · easy

  • A.寄存器← 正确答案
  • B.共享内存
  • C.L2 缓存
  • D.全局内存

答案:A。寄存器位于每个线程“腰间口袋”位置,访问只需约 1 周期;共享内存和 L1 约 20~35 周期,L2 约 200 周期,全局内存高达 400~800 周期。

Q2. 共享内存与 L1 缓存延迟接近,但它与 L1 最大的区别是什么?

单选题 · easy

  • A.由程序员显式管理← 正确答案
  • B.容量大一百倍
  • C.全局统一编址
  • D.硬件自动替换数据

答案:A。L1 缓存由硬件自动管理,程序员无法干预替换策略;共享内存则是程序员在代码里显式分配、显式搬运的“工作台”,可控性更强,也更容易出错。

Q3. 一个 warp 的 32 个线程访问共享内存时,两个线程访问同一 bank 的不同地址,这叫什么?

单选题 · mid

  • A.缓存未命中
  • B.存储体冲突← 正确答案
  • C.内存对齐错误
  • D.合并访问

答案:B。共享内存有 32 个 bank(抽屉),不同线程访问同一 bank 的不同地址时发生“抽屉挤兑”,访问被迫串行。这是共享内存编程特有的性能陷阱,与缓存未命中无关。

Q4. 以下哪些做法能提高全局内存带宽利用率?

多选题 · mid

  • A.让相邻线程访问连续地址← 正确答案
  • B.避免共享内存银行冲突
  • C.按 32 字节扇区对齐访问← 正确答案
  • D.尽量多用寄存器

答案:AC。合并访问的核心就是连续 + 对齐,让每次内存事务满载。银行冲突针对的是共享内存而非全局内存;多用寄存器未必提高全局带宽,反而可能降低占用率。

Q5. 算术强度 I = 运算量/字节数,它主要用于判断程序的什么性质?

单选题 · mid

  • A.显存容量需求
  • B.计算密集还是内存密集← 正确答案
  • C.寄存器分配数量
  • D.缓存命中率高低

答案:B。算术强度描述每搬 1 字节数据做多少次运算。I 小说明访存量大、受带宽限制;I 大说明计算为主、受算力限制。它是 Roofline 模型的核心输入。

Q6. GPU 上每个线程占用过多寄存器,最直接的后果是什么?

单选题 · mid

  • A.共享内存容量变大
  • B.驻留线程数下降← 正确答案
  • C.全局内存容量变小
  • D.L2 缓存命中率提高

答案:B。寄存器文件总量固定,人均占用越多,SM 能同时容纳的线程越少,用线程并行掩盖访存延迟的能力被削弱。这就是为什么“省寄存器”能提升占用率。

Q7. Roofline 模型中,当算术强度 I 小于算力/带宽交界点时,性能瓶颈在哪个环节?

单选题 · hard

  • A.计算单元
  • B.内存带宽← 正确答案
  • C.编译优化
  • D.共享内存容量

答案:B。Roofline 中陡坡段受带宽限制。当 I 小于交界点(ridge point)时,即使算力再强也因数据喂不饱而跑不满,此时优化重点是增大复用、减少搬运。

Q8. 下列哪些属于对“局部性原理”的正确利用?

多选题 · mid

  • A.反复复用同一共享内存块← 正确答案
  • B.随机跳变地址读取
  • C.内层循环遍历连续数组← 正确答案
  • D.每次只读一个新地址

答案:AC。时间局部性对应“反复复用同一块数据”,空间局部性对应“连续访问相邻地址”。随机跳读和每次读新地址都破坏局部性,导致缓存命中率下降。

Q9. GEMM 分块优化主要降低了哪一层次的数据搬运量?

单选题 · hard

  • A.寄存器到共享内存
  • B.全局内存 DRAM 流量← 正确答案
  • C.共享内存到寄存器
  • D.线程束间同步

答案:B。分块把大矩阵切成适合共享内存容纳的小块,使每个数据块被反复复用,从而显著减少全局内存的往返次数。片上各级之间本来就快,重点省的是最慢的 DRAM 搬运。

Q10. 关于 FlashAttention 对显存层次的利用,下列说法正确的是?

多选题 · hard

  • A.将 QKV 分块载入共享内存← 正确答案
  • B.减少与全局内存的数据往返← 正确答案
  • C.不需要再使用全局内存
  • D.通过增加寄存器数量消除注意力计算

答案:AB。FlashAttention 的核心是用共享内存承载 QKV 分块,让注意力计算在片上完成,把全局内存访问量从 O(N²) 降到 O(N)。它并非不用全局内存,也没有靠增加寄存器来消除计算。