GPU流多处理器SM架构

可视化演示

展示GPU SM内部的CUDA Core、Tensor Core协同与Warp调度机制

SM基础架构

想象一家开在繁华商圈的大型连锁餐厅。总部的设计图纸上，每一层楼都被规划成几十个一模一样的小后厨：同样的灶台数量、同样的橱柜、同样的排风扇，配置完全一致。为什么不像单户餐厅那样只做一个气派的大厨房？因为与其费尽心思造一个能同时炒一百道菜的巨型灶台，不如复制一大堆普通厨房让它们同时开工——分工协作，总吞吐量反而更高、也更可靠。

GPU芯片的设计者正是这么想的。GPU内部并没有“一个超级处理器”，而是集成了几十乃至上百个结构相同、相互独立的计算单元块，每个单元块都能独立取指令、独立做调度、独立执行计算。这些单元块就是流多处理器SM。打开一张Ampere架构的芯片白皮书，你会看到占据芯片绝大部分面积的不是某一个庞然大物，而是一排排整齐的SM方块，它们通过片上的互联网络与L2缓存、显存控制器相连。

一个典型的SM内部包含四类关键部件：第一，CUDA核心阵列——几十到上百个能执行整数、浮点运算的通用计算单元；第二，张量核心——专门加速矩阵乘法的硬件单元；第三，存储体系——寄存器文件、共享内存以及各级缓存；第四，控制逻辑——warp调度器和指令分发单元。从Fermi到Ampere再到Hopper，SM内部的布局不断调整，但“多模块协同、各司其职”的设计思想始终没变。GPU的并行度，本质上是SM数量的线性放大：

就像餐厅的繁荣靠的不只是单个厨师的厨艺，更是后厨的复制与排布，GPU的算力也深深扎根于SM的架构设计。理解了这个“后厨单元”，接下来我们就推门进去，看看里面的灶台和师傅们。

计算核心单元

每个后厨里都配有两种“人力”：全能厨师什么菜都能炒，从切丝到爆炒样样在行；自动包饺子机则只干一件事——包饺子，但速度是厨师的几十倍。SM里恰好也有这样两类核心，它们各司其职，配合起来威力惊人。

第一类是我们熟知的CUDA核心，它的本质是一个通用的标量运算单元，负责执行整数和浮点数的加减乘除、逻辑判断、比较以及融合乘加（FMA）等指令。所谓FMA，就是“乘法和加法一步完成”：计算 a = b × c + d 时，普通流程先乘后加两步，FMA用一个单元一口气算完。CUDA核心的优势在于“全能”——任何指令都能处理，哪怕遇到分支跳转也能灵活应对，因此适合控制流复杂、计算形态不规则的代码。

第二类是张量核心（Tensor Core），它是NVIDIA从Volta架构起为深度学习量身定制的“包饺子机”。一个张量核心每个时钟周期就能完成一个4×4×4矩阵的乘加运算，相当于4096次乘加。为什么专门造这种单元？因为神经网络里最耗时的操作——卷积、全连接、注意力——归根结底都是巨型矩阵乘法。把矩阵乘扔给张量核心，把元素级的辅助运算留给CUDA核心，这座后厨的产出效率就能成倍提升：

训练大模型时，绝大部分矩阵乘法都会被调度到“包饺子机”上完成，全能厨师则负责激活函数、归一化、数据搬移等杂活。两类核心的无间配合，才是SM能撑起千亿参数训练的关键。

存储层次

厨师的工作效率不仅取决于手艺，更取决于“取料”的速度。最贴身的是围裙口袋里随时能摸到的盐瓶，取用只需一眨眼；稍远一点是后厨正中央的大理石备菜台，大家手边的半成品都码在上面；最远的则是地下室的冷藏仓库，容量大得惊人，但跑一个来回要好几分钟。SM的存储设计，完全复刻了这套“贴着算、就近放、远处存”的哲学。

最贴身的是寄存器文件（Register File）。它紧贴着CUDA核心，每个线程都有自己私有的几个寄存器，访问延迟只有约1个时钟周期，是所有存储中速度最快的。但寄存器是稀缺资源——A100每个SM的寄存器文件只有256KB，平均下来每个线程大约只能分到几十个寄存器。接下来是共享内存（Shared Memory），它同样位于SM内部，但属于同一个线程块的所有线程共用，大家可以通过它交换数据、协作计算，访问延迟约20~30个周期。共享内存需要程序员在代码里显式声明和搬运数据，相当于“备菜台要自己摆盘”。最外围则是全局内存（显存），容量可达数十GB，但延迟高达400~800个时钟周期——去仓库取一趟菜的时间，够灶台炒几十道菜了。SM内共享内存的带宽估算公式如下：

在CUDA优化中，数据放对位置比什么都重要。把反复用到的数据从仓库搬上备菜台（共享内存），把最热的数据塞进口袋（寄存器），让访问全局内存的次数降到最低——这套“就近取料”的思路，正是FlashAttention等经典算子的核心设计逻辑。

Warp调度机制

主厨不会一盘菜一盘菜地炒，而是把订单按32桌合成一组，整组一起下锅。哪一组食材先到就先做哪一组；某组还在等配菜，主厨就转身去炒别的组。后厨里最忌讳的事是“人等灶”——订单可以等，灶台绝不能闲着。GPU里的调度，玩的就是同一套把戏。

Warp是GPU硬件的最小调度单位：32个线程被捆绑为一组，始终执行同一条指令，这就是SIMT（单指令多线程）模型。当你在CUDA里启动一个包含1024个线程的线程块时，硬件会把它自动切成32个warp，逐个派发执行。每个SM内部有4个独立的warp调度器，每个调度器每个时钟周期可以从“就绪”的warp中挑选一个，向计算核心发射指令。这里的关键在于延迟隐藏：全局内存访问动辄几百个周期的延迟，但这期间调度器并不等待，而是立刻切换到另一个不依赖该数据的warp继续执行——“人歇灶不歇”。SM能同时容纳的warp数量有限（如A100为64个），实际驻留的warp数与上限的比值，就是GPU圈子里常说的占用率：

只要排队等菜的人足够多，主厨就永远有活干；同理，只要驻留的warp足够多，计算单元就不会饿肚子。这也是为什么CUDA优化中，经常用“提高占用率”来对抗访存延迟——不是为了多算，而是为了“总有活可干”。

任务组织模型

餐厅总部接到一个两千人的婚宴订单，当然不会把活儿全压给某一个后厨。销售部会把宴席拆成每桌8~10人的小桌，每张桌子的菜单装进一个信封，再按后厨的忙闲程度分发给各楼层的厨房。CUDA把任务派发给SM的方式，和这套流程如出一辙。

程序员在CUDA中启动内核时，会定义一个“网格”（Grid），网格是“线程块”（Block）的集合；每个线程块包含若干线程，比如128、256或512个。硬件拿到这个网格后，会把一个个线程块按顺序分配给各个SM——每个线程块一旦被分配到某个SM，就会“定居”下来，直到所有线程执行完毕才会腾出位置。线程块内部，线程之间可以通过共享内存交换数据，并借助 __syncthreads() 设置“集合点”：块内所有线程必须都到达这个点，才能继续往下走，就像一桌客人必须全部到齐才能开席。而线程块之间完全独立，互不感知。每个SM能同时驻留多少个线程块，由硬件容量和块大小共同决定：

把订单切成合适大小的“桌”，才能让每个后厨都恰好满负荷，又不至于让一桌客人跨两个厨房。实际工程里，block大小选128还是512，直接影响warp的划分、共享内存的占用和最终的负载均衡——这正是每一位CUDA开发者反复调参的核心原因。

深度学习优化

想让后厨的产出翻倍，聪明的老板不会只给厨师发更多调料，而是重新设计流程：包饺子不用手工改走机器，配菜提前切好码上备菜台，采购一趟就把整周的食材拉回来。深度学习在GPU上的性能优化，本质上就是这三板斧。

第一板斧：把矩阵乘法交给张量核心，并用“分块”（Tiling）策略让数据在共享内存里反复复用。以经典的FlashAttention为例，它不再把巨大的注意力矩阵整体算完再写回显存，而是把Q、K、V切成一块块小片，轮流搬进共享内存，在SM内部完成分块计算后再写回——访存量从O(N²)级别降到了O(N)级别，训练速度成倍提升。第二板斧：让访存“合并”。一个warp里的32个线程如果读写连续的地址，硬件可以一次事务把所有数据取完；如果各自访问分散的地址，就要拆成几十次事务，效率暴跌。第三板斧：用“算术强度”判断该优化什么。这个指标衡量“每读一个字节的数据，能做多少次浮点运算”：

TensorRT、FlashAttention、cuBLAS这些名字之所以如雷贯耳，正是因为它们在帮你的“后厨”重新设计流程：让包饺子机满转、备菜台高效周转、仓库访问次数压到最低。理解了SM的内部结构，你就能看懂这些优化工具的每一招，甚至自己写出下一个FlashAttention。

AI 应用场景

Python 演示

"""SM集群调度模拟：演示Warp轮转调度与延迟隐藏"""from collections import deque MAX_WARPS = 6  # SM 最多容纳的活跃 Warp 数  class Warp:    """一个 Warp 代表 32 个线程构成的指令流"""     def __init__(self, wid, inst_count=6):        self.wid = wid        self.inst_left = inst_count        self.ready_at = 0  # 本次访存返回的全局周期     def issue(self, cycle):        """执行一条指令，返回 True 表示整个 Warp 已完成"""        self.inst_left -= 1        if self.inst_left % 2 == 1:  # 每 2 条指令模拟一次访存等待            self.ready_at = cycle + 3        return self.inst_left <= 0  class SM:    """流多处理器：单个调度器对活跃 Warp 做轮转调度"""     def __init__(self):        self.active = deque()        self.done = []     def add(self, warp):        if len(self.active) < MAX_WARPS:            self.active.append(warp)            return True        return False     def tick(self, cycle):        if not self.active:            return        warp = self.active.popleft()        if warp.ready_at > cycle:            self.active.append(warp)  # 访存未返回，先切到别的 Warp            return        if warp.issue(cycle):            self.done.append(warp.wid)        else:            self.active.append(warp)  def main():    sm = SM()    total = 12    index = 0    cycle = 0    while index < total or sm.active:        while len(sm.active) < MAX_WARPS and index < total:            sm.add(Warp(index))            index += 1        sm.tick(cycle)        if len(sm.done) == total:            break        cycle += 1    print(f"共调度 {total} 个 Warp（每个含 32 线程），耗时 {cycle} 周期")    print("完成顺序:", sm.done)  if __name__ == "__main__":    main()

1. self.active = deque()活跃Warp队列，对应SM调度器手里的“待排班表”。
2. if warp.ready_at > cycle:判断访存是否已返回，未返回就放回队尾，先执行别的Warp。
3. while len(sm.active) < MAX_WARPS and index < total:队列有空位就补充新Warp，对应线程块分批入驻SM的过程。
4. sm.tick(cycle)每个时钟周期调度器执行一次派发，对应硬件每周期发射一条指令。
5. print(f"共调度 {total} 个 Warp...")输出总耗时，直观感受延迟隐藏对吞吐的贡献。

常见误区

面试题