所以当我们聊AI芯片的时候到底在聊什么如果只看厂商发布会上的TOPS、TFLOPS这些数字很容易陷入一种参数拜物教——好像算力堆上去一切问题就都解决了。但做过实际AI项目的人心里都清楚真正决定一个模型训练快不快、推理顺不顺的恰恰是那些发布会PPT上很少展开讲的硬件底层机制。而在这些底层机制里最核心、也最容易被当成理所当然的就是GPU的执行单元。这篇文章我想把GPU执行单元这件事彻底聊透。不仅告诉你它是什么更想拆开给你看它怎么工作、为什么这样设计、以及AI计算和它之间那种近乎天作之合的匹配关系到底建立在什么原理之上。不管你是刚入门CUDA编程的开发者是在做模型训练/推理优化的算法工程师还是纯粹对AI芯片底层架构好奇的读者这篇文章都会给你一个从能用GPU到懂GPU的认知台阶。1. 为什么CPU搞不定的活要专门造一个GPU来干1.1 执行单元的起源从计算到并行计算的岔路口在讨论执行单元之前得先把一个根子上的问题说清楚CPU里也有计算单元为什么AI计算最后还是选择了GPU回到计算机体系结构的基本盘。CPU的设计哲学是低延迟控制流优先它要应对的是指令高度分支、数据依赖复杂、逻辑判断密集的通用任务。为了做到这一点CPU把大量的晶体管预算花在了分支预测器、乱序执行引擎、大容量缓存这些辅助设施上。真正干活的ALU算术逻辑单元反而在芯片面积里占比不高。GPU则走了另一条路。它的假设是如果我面对的工作负载本身就包含海量的、互不依赖的、可以同时执行的计算那我为什么要把晶体管浪费在预测分支和乱序调度上不如把每一分资源都堆到计算单元上让几千个计算单元同时开工。这就是执行单元在GPU上成规模出现的原因——不是为了替代CPU而是为了承接CPU不适合做的重计算、轻控制的工作。1.2 执行单元到底是一块还是一堆必须先把概念理清。当你听到GPU执行单元这个词时要意识到它是一个层级概念。最宏观的层面整颗GPU芯片可以被视作一个由海量执行单元组成的计算阵列。往下一层GPU内部被划分为若干Graphics Processing ClusterGPC每个GPC里有若干Streaming MultiprocessorSM——在AMD平台上叫Compute UnitCU在Intel的Xe架构里叫Execution Unit或Xe Core。这些SM/CU/Execution Unit才是真正意义上的执行单元核心。再往下一层每个SM内部还有数十到上百个更细粒度的处理元件NVIDIA叫它们CUDA Core或者说流处理器AMD叫Stream ProcessorIntel叫Xe Vector Engine。所以执行单元这个话题本质上是一个纵向贯穿整个GPU硬件体系的大话题。不是说SM是执行单元而CUDA Core不是而是它们各自处在执行这个链条的不同层级上各自承担不同的职责。1.3 为什么AI芯片领域如此看重执行单元答案不那么玄妙因为AI计算就是建立在大量简单算术操作并行执行之上的。无论是卷积、全连接、矩阵乘、Attention里的QKV投影还是Transformer里每一步非线性激活拆到最底层就是成千上万亿次的乘法和加法。这些东西没有复杂的控制依赖没有诡异的分支跳转只有一点——量大、管饱。举个例子一个7B参数的LLM做一次前向推理参数量大约70亿哪怕只处理一个token也要把这70亿个参数全部读出来参与计算。对应的乘加操作数轻松突破百亿级别。这种规模的纯算术负载如果靠CPU那十几个核心慢慢跑根本不可能支撑起如今的大模型应用。GPU执行单元就是为了以最大吞吐量完成海量无关算术操作而存在的。理解了这一点后面所有的架构细节都变得顺理成章。2. 一块GPU芯片的计算骨架从GPC到SM再到CUDA Core2.1 NVIDIA Ampere架构下的芯片解剖为了让讨论不那么飘在空中我以目前生产环境中最常见的NVIDIA Ampere架构为例把一颗典型的GPU芯片比如GA102做个自上而下的拆解。一颗GA102芯片总共有7个GPC。每个GPC内部有16个SM部分GPC可能少一两个所以整个芯片一共是84个SM左右。每个SM里有128个CUDA CoreFP32单元、64个SFU特殊函数单元、4个Tensor Core三代Tensor Core每个Tensor Core内部实际上是一个四元组阵列、4个Texture Unit、以及对应的LSU载入/存储单元、寄存器堆Regster File和共享内存Shared Memory。这一层拆解下来你会发现执行单元的含义已经分叉了**如果你把SM看成执行单元那它是一组混合功能单元的组合体如果你把CUDA Core或Tensor Core看成执行单元那它们是能完成特定类型算术运算的最小硬件单元。**这两种理解都对关键看你讨论问题的粒度。2.2 SM内部各功能单元的职责切分在SM内部各功能单元的分工是非常明确的。这里用表格来对比会更直观功能单元主要职责适用的计算类型备注CUDA CoreFP32/INT32单精度浮点和整数运算通用计算、激活函数、非矩阵类标量运算每个SM有128个是绝对的数量主力SFU特殊函数计算倒数、平方根、三角/指数函数比如LayerNorm里的除法和开方softmax里的exp数量远少于CUDA Core通常每SM只有32个调用开销高Tensor Core混合精度矩阵乘加GEMM、卷积、Attention投影、MLP层专为AI设计单位功耗算力远超CUDA CoreLSU从全局内存/共享内存加载数据、写回结果所有需要访存的指令访存指令的吞吐直接决定SM能否吃饱寄存器堆暂存线程的局部变量和中间结果所有指令的数据源头/出口容量有限溢出会导致严重的性能下降共享内存SM内部各线程间数据交换的快速存储数据复用、tile搬运、跨线程通信容量小带宽极高是优化的关键战场从这张表可以看出GPU的SM不只是一群计算单元的简单堆叠而是一个带有完整数据通路和存储体系的微型处理器。计算单元只是这个微型处理器里负责算的那部分它能不能跑满取决于数据能不能按时喂到嘴边——这就是为什么后面要专门讲访存和调度。2.3 CUDA Core和Tensor Core的本质差异很多初学者会把CUDA Core和Tensor Core混为一谈这是执行单元话题里最容易出现的认知混淆。这两者的差异说直白点就是CUDA Core是通用工具什么都能干但干矩阵乘这种大活效率一般。每个CUDA Core本质上是一个支持FMA融合乘加的浮点单元一条指令完成da*bc这个操作吞吐是一拍一个。128个CUDA Core同时工作一个SM一个时钟周期能做128次FMA。Tensor Core是专用工具只会干矩阵乘加但干这个活的速度是CUDA Core的好几倍。它是专门为深度学习设计的硬件单元执行的是DA*BC这种矩阵级别的操作。A、B、C、D都是矩阵不是标量。以Ampere代的三代Tensor Core为例一个Tensor Core一个时钟周期能完成一个4×4×4的矩阵乘加等效于执行了64次FMA。所以一块SM里虽然只有4个Tensor Core但它们加起来的矩阵计算吞吐远超那128个CUDA Core。打个比方CUDA Core是能做各种精细活计的工匠Tensor Core是只做标准门窗的流水线工厂。做特殊造型各种非矩阵计算得靠工匠批量盖楼矩阵运算工厂的效率碾压工匠。3. 指令在SM里到底是怎么跑起来的3.1 Warp执行单元调度的时间切片有了执行单元下一步自然要知道指令怎么在这些单元上跑。NVIDIA的做法是把32个线程绑定为一个Warp以Warp为单位进行调度和执行。这是什么概念就是你写CUDA代码的时候虽然你定义的是成千上万个线程但硬件执行时SM里真正的最小调度单位是Warp。一个Warp里的32个线程执行同一条指令只是各自处理不同的数据。这就是SIMT单指令多线程模型。从SM内部看一个SM通常有4个Instruction Scheduler指令调度器每个调度器每个时钟周期可以向相应的功能单元发出一条指令。但注意这条指令是一个Warp的指令也就是说一条指令发出后对应的功能单元要为这个Warp里的32个线程同时干活。麻烦就出在这里。一个Warp有32个线程要做FMA但每个SM只有128个CUDA Core。调度器一条指令要让32个线程同时执行FMA的话就需要32个CUDA Core。那其余96个Core干什么答案是另外还有三个调度器在同时发出指令各自服务一个Warp。四个调度器一拍发出四条指令每条指令占32个线程的资源总共占满128个CUDA Core。所以**一个SM能同时处于执行状态的Warp数量理论上受限于它有多少个可用的功能单元。**这就是吞吐的本质。3.2 延迟隐藏为什么GPU不怕慢CPU应对访存延迟的方式是缓存命中、乱序执行、分支预测核心思路是别让流水线停下来。GPU的方法则完全是另一个次元打不过就用车轮战。GPU有一个术语叫Occupancy占用率指的是SM上同时活跃的Warp数量。为什么GPU要承载这么多Warp就是为了延迟隐藏。当一个Warp的指令需要从全局内存取数据可能需要几百个时钟周期才能拿到结果。这段时间里如果调度器等这个Warp那执行单元就闲置了性能灾难。但如果SM上有足够多的Warp调度器就可以立刻切到另一个已经准备好数据的Warp让执行单元永远有活干。这就是为什么你在编程时要关注有没有足够多的线程/Block来填满SM。线程太少执行单元就会因为等待Data而挨饿线程太多寄存器不够用又会导致本地内存溢出反而拖慢速度。找到这个平衡点是每个CUDA性能优化工程师的基本功。3.3 分支分化执行单元最怕的程序结构SIMT的执行方式带来一个经典问题分支发散或叫分支分化。想象一下一个Warp里的32个线程执行if (data[i] 0) { do_a(); } else { do_b(); }。如果这32个线程的数据有的大于0有的小于0那么执行单元没法同时执行do_a和do_b——它只能先让走true分支的线程执行do_a此时走false分支的线程只能在一旁干等被掩码屏蔽然后再反过来执行do_b。这个Warp一共执行了do_a和do_b两段代码耗时翻倍但有效计算量只有一半。这直接导致了GPU执行单元的利用率下降。所以我们在写kernel时要想尽办法避免Warp内分支可以把数据预先排序、把不同分支拆成独立的kernel、或者用算术方法替代分支。这些都是实际项目中能实打实省出几十毫秒的优化手段。3.4 访存合并让LSU吃饱的秘诀除了分支GPU执行单元最在意的事情就是访存模式。LSU载入/存储单元负责把数据从显存搬进寄存器或者把运算结果写回显存。当一个Warp的32个线程访问全局内存时如果它们的地址是连续的比如访问数组的相邻元素硬件可以把这个访问合并成少数几次大的内存事务——这种模式叫合并访存coalesced access。反过来如果32个线程各自访问毫无规律的地址那内存控制器就要拆成几十次小事务效率一落千丈。这个特性和执行单元的关系太直接了**访存效率低再快的执行单元也得空转等数据。**很多人在写kernel时觉得自己的计算逻辑没问题跑出来却慢得像爬一查profile大概率是访存没合并。4. AI计算如何与执行单元互相成就4.1 一个GEMM在SM层面是怎么被切分的深度学习里最核心的运算GEMM通用矩阵乘是理解执行单元价值的最佳窗口。假设要计算C[M,N] A[M,K] * B[K,N]M、N、K都是几千的规模。GPU的做法不是把一个巨大的矩阵整体丢给执行单元那样数据放不下计算也组织不起来。实际的策略是分块tiling在Block层面把最终结果矩阵C划分成若干个Block Tile比如每个Block负责一个128×128的C子块。在Warp层面一个Block内部又由多个Warp各自负责其中一部分比如每个Warp处理一个64×64的子块。在指令层面每个Warp的线程调用mma.sync指令Tensor Core指令一次完成m16n8k16之类的矩阵分块乘加。这个分层切分策略本质上就是为了让数据尽可能留在SM内部的共享内存/寄存器里减少对全局内存的访问。每个Thread从全局内存中只读取自己负责的那一小块A和B数据存入共享内存然后反复复用——这正是计算密度的核心执行单元在一个数据上做的计算次数越多访存带宽压力就越小。4.2 Tensor Core里的乘法累加洪流Tensor Core执行GEMM的微观机制非常有意思。以A100的第三代Tensor Core为例它执行的是D A * B C其中A、B、C、D都是4×4的矩阵。也就是说一次Tensor Core指令可以完成4行4列的A矩阵16个FP16元素4行4列的B矩阵16个FP16元素4行4列的C矩阵16个FP32元素输出4行4列的D矩阵16个FP32元素等效于64次FMA运算。而且Tensor Core支持在一个时钟周期内完成这个操作。对比一下一个CUDA Core一个周期只能完成1次FMA。虽然Tensor Core占用的晶体管面积比CUDA Core大不少但就矩阵运算而言它的能效比是碾压级别的。后续代际的Tensor Core规格继续膨胀H100的第四代Tensor Core直接支持m16n8k4到m16n8k16等多种形状还有FP8支持到了H200和B200时代Tensor Core的矩阵形状进一步扩大单指令的FLOPs还在攀升。可以说大模型的每一次规模跃迁背后都是Tensor Core这个专用执行单元的算力密度升级。4.3 训练和推理对执行单元的需求为什么不同聊执行单元必须区分训练和推理两种场景因为它们对硬件的压力点完全不一样。训练阶段的特点是海量计算、数据精度要求高、需要反向传播求梯度。此时执行单元的压力是持续满载GPU需要以最高吞吐执行前向和反向的GEMM而且通常使用FP16/BF16精度配合混合精度训练因为梯度更新和损失计算需要FP32精度。训练时执行单元最容易被算力瓶颈卡住因为计算量实在太大每个时钟周期都要最大程度地压榨Tensor Core。推理阶段则复杂很多。在线推理通常有延迟要求Batch Size往往很小甚至为1。此时执行单元可能吃不饱因为单次请求的计算量太小内存带宽反而成了瓶颈。这就是为什么很多推理优化工作会做Continuous Batching连续批处理、PagedAttention本质上都是想办法把不同的请求拼成更大的矩阵让执行单元不至于因为活太少而空转。另外推理还常常用量化INT8/FP8把执行单元的单位算力进一步转化成吞吐优势。4.4 Roofline模型判断程序瓶颈的一把尺理解了执行单元和访存系统你就有了一把分析程序瓶颈的尺子Roofline模型。Roofline的核心思想是一个kernel能达到的实际性能不大于算力上限和带宽上限两者中的较小者。如果一个程序的计算密度每字节访存对应的计算量很低那么它大概率是访存瓶颈反之如果计算密度很高但执行单元已经被压满那就是算力瓶颈。这张屋顶图像是在如何在GEMM优化、卷积算子开发时做决策的关键依据如果一个op是访存瓶颈那你去调Tensor Core的指令形状一点用都没有应该去优化共享内存的复用、改善访存合并如果一个op是算力瓶颈那就要考虑能否用Tensor Core替代CUDA Core、能否用更低精度的数据格式增加单指令的FLOPs。5. 谁在喂执行单元存储与调度体系5.1 寄存器堆、共享内存与全局内存的分工执行单元本身跑得快但数据链路跟不上一切白搭。GPU里存在一个金字塔形的存储体系寄存器堆每个SM里有256KB的寄存器堆Ampere分割给所有活跃线程。这是速度最快的存储执行单元做任何计算数据基本都要先在寄存器里。但寄存器是隐私的一个线程的寄存器其它线程访问不了。共享内存每个SM有最多164KB可配置的共享内存整个Block的所有线程都能访问。这是执行单元之间协作的最佳场所也是tiling策略发挥作用的地方。共享内存速度远快于全局内存相当于一个用软件管理的二级缓存。全局内存就是显存HBM/GDDR所有Block都能访问容量大但延迟高。执行单元从全局内存取数据是最贵的操作要尽量避免。理解了这套体系你就知道为什么优化CUDA kernel的核心思路永远是别让执行单元频繁去全局内存取数把数据尽量往共享内存和寄存器里搬。5.2 调度器如何决定下一个执行谁前面提过SM里有多个Instruction Scheduler那调度器到底按什么规则选WarpNVIDIA的Warp调度策略在不同代际略有差异大体上是一个循环式的优先级调度——调度器维护一个Warp列表每次选一个准备好了的即依赖数据已到齐、且没有被阻塞Warp发出下一条指令。这个选择过程每个时钟周期都在进行目的是最大化功能单元的利用率。有一个值得了解的概念是两个Warp的发射间隔issue interval。如果两条指令之间没有数据依赖那它们可以流水线式地连续发射如果有依赖就必须等前一条指令的结果写回。这要求开发者对指令级并行有意识在写代码时尽量让计算链不要过长。当然对绝大多数应用开发者来说这个问题已经被编译器处理掉了大半但知道底层逻辑还是有帮助的。5.3 为什么说SM的占用率不是越高越好这是个反直觉但极其重要的优化细节。Occupancy高意味着SM上有更多Warp可以切换延迟隐藏效果更好。但代价是每个Warp分到的资源变少——尤其是寄存器。如果你的kernel因为占用率太高一拍每个线程只分到少量的寄存器被迫把局部变量溢出到本地内存其实落在L1/全局内存里那性能反而会暴跌。在实际调优中我经常用__launch_bounds__来限制一个Block里的线程数从而给每个线程争取更多寄存器换来更高的单线程指令执行效率有时性能反而比强行拉满Occupancy更好。这就是执行单元话题里最有意思的平衡术你要的是效率而不是占用率数字。6. 从CUDA代码到执行单元一次性能优化的实战视角6.1 一个矩阵乘kernel的优化路径纸上谈兵够了用一个实战视角来看执行单元是如何被优化的。假设我们要实现一个FP16 GEMM的CUTLASS式kernel优化的路径大致是朴素版本每个线程直接算C中的一行一列访存完全不合并执行单元利用率低得可怜。共享内存分块把A和B矩阵分块载入共享内存每块被多个Warp复用大幅减少全局内存访问。这个阶段执行单元开始能吃上热饭。寄存器优化每个线程用寄存器数组手动缓存A、B的切片减少重复读共享内存。向量化访存用float4/half2这类向量类型一次读多个float增加访存吞吐。Tensor Core把核心计算换成mma.sync指令将执行单元从CUDA Core切换到Tensor Core算力瞬间翻几倍。Pipeline用CUDA Graph或异步拷贝cp.async实现共享内存的prefetch把计算和访存重叠起来让执行单元永远有数据可算。每一步的改进本质上都是在调整执行单元与数据之间的距离。距离越短性能越高。6.2 用Profile工具看到执行单元的饥饿理论说了一大堆真正要定位性能问题最终还是得靠工具。我推荐从NVIDIA Nsight Compute入手因为它是可以直接观测SM内部执行单元状态的神器。Nsight Compute里的几个关键指标直接对执行单元负责SM BusySM内计算单元活跃占比。如果很低说明执行单元在等数据优先排查访存。Issue Active调度器发出指令的活跃度。如果低说明Warp处于阻塞状态等待数据或同步。Executed Ipc Active平均每周期执行的指令数。接近4Ampere SM有4个调度器说明执行单元喂得很饱远低于4则说明有阻塞。Memory Throughput全局内存吞吐。如果这个指标先打满那就别折腾执行单元了它在等着吃饭。我遇到过很多次代码写得像优化过但实际还是慢的kernel一测Executed Ipc Active只有0.3。这意味着执行单元97%时间在发呆。这种问题你用多少算力纸面数据都救不回来必须从数据流入手。6.3 大模型时代的执行单元真相算力规格之外更值得关注的指标大模型火起来之后各家芯片厂商的算力数字一个比一个吓人。但我要泼一盆冷水在执行单元的话题里峰值算力只是入场券持续稳定的有效算力才是胜负手。有效算力峰值算力×实际利用率。实际利用率取决于你的算子库有没有针对具体硬件做深度优化FlashAttention就是一个绝佳案例取决于你的显存带宽能不能喂饱执行单元取决于你的矩阵形状填充度够不够高比如seq_len3的短序列推理Tensor Core根本跑不满……所以做AI基础设施选型时别只看TFLOPS要看基于你真实模型benchmark出来的有效吞吐。这也是为什么H100、A100这种经过市场反复打磨的芯片在生产环境里地位稳固——不只是因为纸面强是因为在真实负载下它们的执行单元确实能被高比例地喂饱。7. 执行单元设计的边界与未来功耗墙、存储墙与不断扩张的AI需求7.1 功耗墙芯片能跑多快其实由散热决定GPU执行单元的发展视觉上看起来是核心越来越多、频率越来越高但物理上其实一直被功耗墙卡着。制造工艺不断缩小晶体管密度不断提高但单位功耗的散热量并没有跟上传导速度的提升。一块700W功耗的GPU把几乎所有功耗都喂给了执行单元然而这带来的热量极其恐怖。于是我们看到一个趋势算力增长不再主要依靠堆核心数量而是依靠每瓦性能的提升——例如引入MXFP4格式、用更小的数据位宽换更快速度、或者用更微妙的激活稀疏化。执行单元的设计越来越受到能效比约束。谁能在同等功耗下跑出更高的有效算力谁就在AI芯片竞争中占据优势。7.2 存储墙算得再快数据喂不进来也是白搭与功耗墙并行的还有存储墙。HBM的带宽这两年提升很快但相比执行单元的算力提升还是慢了不少。以H100为例FP16稠密算力接近1000 TFLOPS而HBM3的带宽只有3.35TB/s。你可以算一笔账每个字节的数据执行单元可以做约300次FP16运算。如果你的算法不注重数据复用那么即便执行单元全速运转也会被内存带宽卡死。这也是为什么AI推理和训练都要强调计算密度和kernel融合——本质都是为了减少对存储带宽的需求让执行单元在数据复用上多做文章。7.3 下一代执行单元往哪里走展望未来GPU执行单元的设计方向其实已经非常清晰更宽的Tensor Core指令单指令覆盖更大的矩阵形状减少指令发射开销。更低精度的数据格式FP8、MXFP4、甚至更低位的三值/二值量化让执行单元每个周期能做更多次有效乘法。更为异构的SMCPU、GPU、张量单元、稀疏加速器组合共存各司其职。英伟达的Grace Hopper以及更近的Blackwell架构已经在走向把整个计算系统的各个部分都做成专用执行器的路线。chiplet化设计把执行单元按区块拆分到不同小芯片上再用高速互连连接突破单芯片光罩面积的限制。执行单元这个概念的边界会从GPU内部的硬件单元逐渐扩展到整个计算系统的执行能力。但对从事AI开发的人而言最核心的那句话依然没变让执行单元吃饱、吃好、别乱跑。我在实际项目中看到过太多类似的案例一个模型推理慢得离谱开发者第一反应是GPU太弱了要换卡结果一分析发现是kernel写得稀烂执行单元利用率不到20%。换卡之前先做一次Profile通常能分分钟找出问题——无论是warp分化、访存不合并还是Tensor Core压根没用上。这是我在这个领域反复体验到的真理硬件在进步但能否吃透这些执行单元的脾性才是决定你是NVIDIA显卡用户还是GPU性能优化工程师的分水岭。等你真正把一个kernel从能跑优化到跑满看着Nsight Compute里Executed Ipc Active从0.3爬到3.7的时候那种对整个计算体系的理解比任何参数表上的数字都有说服力。
