极客工场

Appearance

共 5,067 字预计阅读约 17 分钟

SM 流多处理器：warp 调度器、寄存器堆、CUDA Core 与 Tensor Core 的内部结构 ​

更新时间：2026-09-25。前置阅读：GPU 架构：warp 调度器 与 CUDA 线程层次。前面讲 GPU 时总说「warp 在 SM 上被调度」，但 SM 不是一个抽象的「核」——它是一块面积可观、内部塞了调度器、上百个执行单元和数百 KB 片上存储的完整子系统。本篇把 SM 这块芯片掰开，看一条 warp 指令从被选中到算出结果，一路上经过了哪些部件。

本文要回答的问题 ​

  • 一颗 GPU 与 SM 是什么关系，为什么说 GPU 性能主要靠「复制 SM」堆出来？
  • 一个 SM 内部有哪些功能单元，各自的典型数量和分工是什么？
  • warp 调度器每周期如何挑选 warp、把指令发射到不同的执行端口，什么是双发射？
  • 寄存器堆为什么要做到 256KB，它和 CUDA Core、shared memory 怎么连？
  • Tensor Core、SFU、load/store 单元这些「专用通道」分别处理什么指令？

一、GPU 就是一堆 SM 的复制：从全卡视角看定位 ​

GPU 的设计哲学是用面积换吞吐：把一个通用的计算单元设计到足够高效，然后在一颗芯片上复制几十上百份，再用片上网络和共享的 L2、显存控制器把它们连起来。这个被复制的基本单元，在 NVIDIA 体系里叫 SM（Streaming Multiprocessor，流多处理器），AMD 叫 CU（Compute Unit），Intel 新架构里叫 Xe Core。名字不同，内部构成相似。

以当代数据中心卡为参照：H100 一颗芯片上有 132 个 SM（完整规格 144 个），A100 有 108 个；消费级 GPU 通常是几十个。每个 SM 独立驻留一批线程块（block / CTA）、独立调度 warp、独立拥有自己的寄存器堆和片上 shared memory。block 之间不共享 SM 内资源——一个 block 一旦被分配到某个 SM，整个生命周期都在那里度过，直到执行完毕，SM 才会腾出位置接收新的 block。

顺带一提 AMD 一侧的对应关系：CU 里调度单位叫 wavefront，宽度是 64 个线程（NVIDIA warp 是 32），CU 内部同样是调度器加多个 SIMD 通道加寄存器堆加 local data share（对应 shared memory）的组合。理解了 NVIDIA 的 SM，再读 AMD 的 ROCm 文档几乎可以逐部件对照。

理解 SM 的关键地位，可以用一句话：线程是程序员看到的抽象，warp 是调度单位，而 SM 是这一切真正发生的物理场所。 显存、L2、NoC 负责把数据送到 SM 门口，而占芯片面积大头的 SM 阵列负责算。所以评估一颗 GPU 的算力，本质上是在数「有多少个 SM、每个 SM 每周期能做多少事、跑多高频率」三件事的乘积。厂商在芯片内部还会把相邻的十几个 SM 编成一个 GPC（图形处理簇），共享几何/光栅化前端，但对计算 kernel 来说这个分组基本透明。本篇聚焦后两个问题：单个 SM 里面长什么样。

二、SM 内部全景：调度区、寄存器堆、执行单元群、片上存储 ​

一个现代 SM 内部分成四块处理区（processing block），每块自带一个 warp 调度器和一套执行单元；寄存器堆、shared memory/L1 则是全 SM 共享的大块 SRAM。下图是 SM 的内部结构骨架。

chip-design/intermediate/sm-cu-internal-structure.md D2 示意图

读图的顺序自上而下：调度器决定「谁、发哪条指令」，寄存器堆提供操作数，指令按类型进入不同的执行单元群，访存类指令最终落到 shared memory/L1 或往外走 L2、显存。下面这张表给出 H100 一档 SM 的典型数量级——数字随架构代际变化（A100 的 FP32 核数是 64），但单元种类和分工是稳定的。

单元典型数量（每 SM）职责
warp 调度器 / 分发单元4 个（对应 4 个处理区）每周期从就绪 warp 中挑选并发射一条指令
寄存器堆1 大块，256 KB（65536 个 32 位寄存器）存放所有驻留线程的私有操作数
FP32 CUDA Core128 个单精度浮点加/乘/融合乘加（FFMA），主力算力
INT32 单元约 64 个整数运算、寻址、比较，可与 FP32 并行执行
FP64 单元64 个双精度浮点科学计算（消费卡常被削减）
Tensor Core4 个一次完成小块矩阵乘加（MMA），AI 算力主力
SFU 特殊函数单元约 16 个sin/cos、倒数、平方根倒数等超越函数
load/store 单元约 32 个量级计算访存地址、发起/完成 load 与 store 事务
shared memory / L1228 KB 合一存储block 内共享数据与一级缓存，切分比例可配

这张表里最值得注意的是「异构」二字：SM 不是 128 个完全相同的通用核，而是一组数量、吞吐、延迟各不相同的专用端口。一条 warp 指令能去哪里，在译码阶段就决定了。

再用「处理区」的视角把数字核对一遍，会更好记：每个处理区拥有 1 个 warp 调度器、32 个 FP32 Core、16 个 FP64 单元、1 个 Tensor Core，外加一份就近的寄存器端口和若干 SFU、load/store 通道；4 个区拼起来就是上表的全 SM 数字。A100 时代每个区只有 16 个 FP32，全 SM 合计 64——两代架构的差距不是换了核心设计，而是同样的结构复制宽度翻倍。这种「小单元 × 多份复制」的组织方式贯穿 GPU 的每一个层级，从全卡复制 SM，到 SM 内复制处理区，如出一辙。

三、warp 调度器：每周期选谁、往哪发、什么是双发射 ​

SM 内的 4 个 warp 调度器各自负责一个处理区。每个调度器手里维护着分派到本区的若干 warp 的状态：程序计数器、就绪位、等待原因。每周期它做两件事——先从就绪 warp（操作数齐备、没有在前序指令上阻塞）中按策略选一个，再把它的下一条指令译码，发射到对应的执行端口。

发射路径是按指令类型分流的：

chip-design/intermediate/sm-cu-internal-structure.md D2 示意图

这张发射图解释了 GPU 延迟隐藏的微观机制：一条 load 指令发射后，该 warp 要等几十到几百周期才有数据，调度器下一个周期不会等它，而是直接选另一个就绪 warp。只要驻留的 warp 足够多（H100 每个 SM 最多驻留 64 个 warp、2048 个线程），执行端口就几乎不会空转。零成本切换的前提是寄存器堆里同时保留着所有驻留 warp 的现场——这正是 SM 需要 256KB 寄存器的根本原因。

调度器判断 warp 是否就绪，依靠一张记分表（scoreboard）：每条在途指令占一个表项，记录它写回哪个寄存器、何时完成；warp 的下一条指令只要源操作数没被在途指令占用、所需执行端口此刻可用，就算就绪。反过来，warp 的挂起原因通常是三类——等显存数据、等 shared memory 事务、等前一条长延迟指令（如 SFU）写回。调度器不需要预测分支，也不需要乱序执行，它的全部智慧都用在「从一批确定就绪的 warp 里公平地挑一个」上。这种用大量独立现场换选择空间的做法，比 CPU 用复杂推测执行隐藏延迟省面积得多。

关于双发射（dual issue），需要区分架构代际。早期 Fermi 架构的一个调度器配两个分发单元，同一条 warp 上两条相互独立的指令（比如一条浮点加加一条访存）可以在同一周期一起发射；Kepler 时代的 SMX 有 4 个调度器，每个也支持双发射，峰值每周期发射 8 条 warp 指令。而从 Maxwell 到当代架构，主流设计改为每个调度器每周期发射一条指令、靠增加调度器数量和执行端口宽度拿吞吐；但「不同类型端口并行」始终存在——同一周期里 4 个调度器可以分别给 FP32、INT32、Tensor Core、load/store 各发一条，这些指令互不冲突。挑选策略大体是就绪 warp 间的轮转加权，硬件不做复杂的优先级排序，但会保证长期公平——一个 warp 不会因为「运气差」长期饿死。对程序员来说，这意味着发射顺序基本不可控，任何两条指令之间都不能隐含「另一个 warp 已经跑完」的假设，跨线程协作只能靠显式同步。分支发散时 warp 内部还需要按路径轮流发射，详见 SIMT 与分支发散。

四、寄存器堆：256KB、全 SM 最大的一块 SRAM ​

寄存器堆是 SM 里最反直觉的部件：习惯了 CPU「每核几十个物理寄存器」的人，很难想象一个 SM 配 65536 个 32 位寄存器（256KB）。但这笔账无法省——一个 SM 上同时驻留最多 2048 个线程，每个线程最多 255 个寄存器，所有线程的现场必须随时可用，warp 切换才能做到零开销。

物理上，寄存器堆并不是一整块大 SRAM 加一个巨型交叉开关，那样布线面积和功耗都不可接受。它按处理区切分并多 bank 组织，每个处理区就近访问自己那一份；一条 warp 指令发射时，由操作数收集逻辑把 32 个线程的源操作数在一两个周期内凑齐，再送进 32 宽的执行通道。浮点中间结果、循环下标、函数参数全部优先落在寄存器里——它是唯一单周期可用、带宽最高的存储层。

还可以从 warp 的角度算一笔现场账：一个 warp 有 32 个线程，若每线程用 64 个寄存器，一个 warp 就占 2048 个；64 个驻留 warp 满编时，接近 13 万个寄存器位槽的需求由这 65536 个物理寄存器按实际占用动态切分——注意上限不是所有线程都能用满 255 个，65536 除以 2048 个满编线程，平均每线程只有 32 个。要让每线程拿更多寄存器，就必须减少驻留线程数，这正是 occupancy 取舍的数学根源。

寄存器堆是共享而有限的资源：SM 上驻留的 warp 越多，每个线程平均分到的寄存器越少。编译器负责分配，若 kernel 每线程寄存器用量过高，能同时驻留的 block 数就下降，调度器手里可换的 warp 变少，长延迟藏不住；再不够用就会 spill 到名为 local memory、实际在显存上的区域，一次访问几百周期。寄存器与 occupancy 的定量关系见 occupancy 与寄存器压力。

五、执行单元群：CUDA Core、SFU、Tensor Core 各司其职 ​

CUDA Core 是最常被提及的单元，但它本质上就是一条流水化的标量浮点（或整数）数据通路：每周期接受一个线程的一对操作数，产出一个结果。SM 标称的「128 个 FP32 Core」意味着一条 FP32 warp 指令（覆盖 32 个线程）由 32 个 Core 一组、在调度区的安排下完成，每个调度区每周期消化一条 warp 指令，4 个区合起来构成 SM 的标称吞吐。INT32 单元独立存在，地址计算、循环控制可以与浮点计算并行，互不抢端口；FP64 单元数量为 FP32 的一半，一条双精度 warp 指令要占两个周期。

**SFU（Special Function Unit，特殊函数单元）**处理不适合通用 ALU 的数学指令：正弦余弦、倒数、平方根倒数、指数对数等。SFU 数量少（约 16 个量级）且吞吐低于 FP32 通道，一条 warp 的 __sinf 往往要多个周期才能退完。密集使用超越函数的 kernel 会明显变慢，工程上常用查表加插值、多项式逼近等方式把热路径上的超越函数换掉。还要注意 SFU 走的是专用近似实现，结果与标准数学库可能有末位差异，对数值精度敏感的迭代算法（如反复收敛的递推）不要想当然地把它和软件实现混用，编译选项里也可以强制用软浮点版本换取位级一致。SFU 的思路与 CORDIC 旋转器用专用结构高效实现函数计算一脉相承。

从吞吐数字上能直观感受各端口的差异：128 个 FP32 Core 满打 FFMA（融合乘加，每周期两个浮点操作）时，SM 峰值是每周期 256 个浮点操作；而 4 个 Tensor Core 执行 FP16/BF16 矩阵指令时，每周期的等效操作数是数千量级——相差超过一个数量级。这也解释了架构文档里一颗 GPU 常常同时标注两个算力数字：通用浮点算力看 CUDA Core，AI 算力看 Tensor Core，二者衡量的根本不是同一种工作负载。

Tensor Core 是 SM 里粒度最粗、算力密度最高的单元。它执行的不是单个数的运算，而是一条矩阵乘加指令：取一小块矩阵（如 16×16 与 16×8，尺寸随代际增大），一个周期级的操作内完成整块乘加。4 个 Tensor Core 的吞吐远高于 128 个 FP32 Core 的总和——这就是为什么矩阵乘法、深度学习训练推理的峰值算力由 Tensor Core 定义，而 FP32 只反映通用计算能力。其工作原理与编程要点见 Tensor Core 矩阵加速，脉动式数据流的背景可读 AI 加速器脉动阵列。

六、load/store 单元与 shared memory/L1：SM 的片上数据门户 ​

执行单元只管算，数据的进出由 load/store 单元负责。它们接收 warp 的访存指令，完成地址计算（基址加各线程偏移）、判断访问落在哪个层级，然后发起事务：命中 shared memory/L1 的在片上解决，超出范围的请求经 NoC 送往 L2 和 HBM/GDDR。一个 warp 的 32 个线程若地址连续，load/store 单元会把它们合并成一个宽事务，这是显存带宽利用率的命门。

顺带一提，L1 这块自动缓存还承接两类容易被忽略的流量：寄存器 spill 的 local memory 访问，以及未显式使用 shared memory 时普通全局变量的命中尝试。所以 shared 配多了、L1 被压小，受影响的不只是普通访存命中率，连 spill 路径的代价都会跟着变差。

shared memory / L1 是 SM 内仅次于寄存器堆的第二大块 SRAM。H100 上合计 228KB，物理上是同一块存储，按配置切出 shared memory（软件显式管理、block 内共享）和 L1 data cache（硬件自动缓存）两部分；A100 上为 192KB，shared 最多约 164KB。shared memory 内部按 32 个 bank 组织，一个 warp 的访问若均匀落在不同 bank 上可一周期并行完成，撞 bank 则串行化。它的完整定位（容量、延迟、与寄存器和 HBM 的比值）在 GPU 内存层级里展开。

store 指令与 load 略有不同：它发射后数据只要离开寄存器、进入写路径，warp 原则上就不必原地死等，写合并与写回由后续流水完成；但 store 之后紧跟同步指令（如 __syncthreads() 或内存栅栏）会强制把这些在途写入排干，这也是栅栏代价的一部分。此外，load/store 单元的地址生成本身也要消耗 INT 通路资源，这就是为什么复杂的下标表达式（多维数组、步长寻址）会让访存密集的 kernel 额外吃整数吞吐。

把 SM 的完整闭环串一遍：block 被分配到 SM，warp 开始执行；调度器每周期从就绪 warp 中选人，寄存器堆提供操作数，指令按类型进 FP32/INT/FP64、Tensor Core、SFU 或 load/store；访存由近及远依次找 shared/L1、L2、显存；遇到长延迟就切换 warp，直到数据回来。单个 warp 看起来走走停停，整个 SM 的各执行端口却被不同 warp 填得满满当当——这就是 SM 的全部设计意图。

这一篇在知识体系里的位置 ​

本篇是 GPU 硬件与编程模型的第三篇：上承 GPU 架构：warp 调度器（全卡视角与延迟隐藏）和 CUDA 线程层次（grid/block/thread 如何映射到 SM），与 SIMT 与分支发散（warp 内指令如何被掩码执行）、occupancy 与寄存器压力（寄存器堆怎么决定驻留 warp 数）互为表里。存储方向接 GPU 内存层级；专用算力深入 Tensor Core 矩阵加速（expert） 与 AI 加速器脉动阵列（expert）。

总结 ​

SM 流多处理器（AMD 称 CU）是 GPU 被复制几十上百份的基本算力单元：一个现代 SM 内含 4 个 warp 调度器，每周期各选一个就绪 warp、把指令按类型分发到 128 个 FP32 CUDA Core（另有约 64 个 INT32、64 个 FP64 单元）、4 个 Tensor Core、约 16 个 SFU 和数十个 load/store 单元；支撑零成本 warp 切换的是全 SM 共享的 256KB 寄存器堆（65536 个 32 位寄存器，按处理区分 bank 切分），而 block 共享数据与一级缓存由 shared memory/L1 这块约 192~228KB 的片上 SRAM 承接；双发射是早期架构（Fermi/Kepler）同 warp 两条独立指令并发的能力，当代设计改为多调度器加异构端口并行——读 SM 的正确姿势不是记数字，而是理解「调度器选人、寄存器出操作数、指令按类型走专用端口、靠大量驻留 warp 把每条长延迟藏进别的 warp 的执行时间里」这条完整闭环。