很多做AI Infra和算力运营的朋友应该都碰到过这种怪事租来的GPU明明标称几百T算力跑大模型推理时nvidia-smi里看到的利用率却只有百分之二三十跑矩阵乘benchmark又能跑出非常漂亮的数字。差距到底出在哪我刚接手算力平台那会儿也被这个问题折磨过后来把底层性能分析彻底摸了一遍才发现大部分人忽略了一个微架构层面的关键角色——指令发射单元。它是AI芯片的“十字路口”每条指令从取指译码后都必须经过这里才能进入真正干活的执行单元。这个路口的调度能力、等待队列和拥堵状况决定了纸面上的算力指标能转化出多少真实吞吐也就是我们常说的算力天花板。这篇文章适合谁看一方面是在做芯片选型、算力平台管理、大模型推理优化的工程师另一方面是纯粹想搞明白“为什么我的卡利用率上不去”的算法同学。我会从微架构原理讲起穿插我实际在GPU和各类AI加速卡上做profile的经验最后落到多卡集群管理层面尽量让不同背景的人都能照着用。1. 标称算力只是理论值指令发射效率才是真实吞吐的阀门1.1 一条指令从取指到执行要过几道门我们先从最基础的过程看起。一段神经网络模型经过编译器最终会变成一连串机器指令。这些指令要先被取指单元从一级指令缓存里一条条取出来送到译码器。译码器做的事情是把操作码、寄存器地址、立即数拆开搞清楚这条指令到底想干什么、要读哪些寄存器、要写到哪去。到这一步为止很多人的理解就结束了总觉得编译完的指令会顺理成章地执行。实际上不是。译码完成之后的指令必须进入指令发射单元由它来判定这条指令能不能立刻往下走。要检查的东西至少有三项操作数是不是已经就绪——比如上一条指令刚说要写某个寄存器结果还没写回来这条指令就去读那肯定不行执行单元是不是有空闲——矩阵单元还在跑上一条大矩阵乘新的矩阵指令就只能等着寄存器和访存端口有没有冲突——有的芯片寄存器堆读写端口有限两条指令同时要读同一组端口总得有个先来后到。全部通过之后指令才被真正“发射”到执行单元上。这个过程用交通来类比特别贴切。执行单元是城市里的各条车道指令发射单元就是进车道之前的十字路口。路口每个周期能放行多少辆车决定了后面车道修得再宽、跑车性能再好都没用。很多时候我们嫌某个AI芯片“跑不出标称算力”本质不是后面的MAC阵列不够快而是这个路口在过去车。1.2 IPC、发射宽度和发射槽利用率衡量路口通畅程度的三把尺子要量化这个路口的通行能力有几个基本指标得认识。第一个是IPC平均每个时钟周期完成或写回的指令数。老玩家看CPU评测时总能看到这个数字它反映的是整个从取指到执行的回流是否顺畅。IPC过低说明指令流里存在大量等待发射端和计算端没有咬合好。第二个是发射宽度也就是每个时钟周期最多能发射几条指令。这里的“条”在不同架构里含义不同。在GPU里指一个调度器每周期可以发射一条warp指令一条warp指令对应32条线程的同一操作在VLIW风格的NPU里指一条超长指令字里有几个作用域可以同时触发比如同时控制矩阵单元、向量单元和DMA引擎。第三个也是我平时最关注的发射槽利用率。它等于实际发出的指令数除以理论可发射的槽位数。很多kernel跑下来这个值可能只有40%到50%意味着每两个周期里有一个周期的发射机会是彻彻底底空着的。GPU的profiler里有个专门的统计项叫Issue Active说的就是这个比例还有一个对应的No Eligible表示调度器想发射但找了一圈发现所有warp都不满足发射条件只能空转。我自己的经验是这三把尺子得配合着看。SM的占用率很高、但Issue Active很低说明执行单元虽然都在线发射端根本没有把计算指令喂进来大概率是访存或者依赖把warp卡住了如果SM占用率也很低那就可能是kernel启动不足比如batch太小整个SM里的线程块没排满。1.3 为什么标称TFLOPS只是理论上限标称算力这个概念本身有个隐含假设每个时钟周期芯片上所有的矩阵单元都在满负荷跑矩阵运算。这个假设要成立必须满足两个条件。第一矩阵指令是连续不断送到MAC阵列的中间不能有断流第二每条矩阵指令要读的数据已经全部准备好放到共享内存或者寄存器里了。第一个条件测的就是指令发射单元。只要发射断流不管标称TFLOPS是300还是1000MAC阵列都只能原地等待。第二个条件其实也间接和发射有关——如果数据搬运指令没有提前发射出去后面的计算指令即便已经轮到发射了也会因为操作数没就绪而发不出去。举个例子A100的SM里有4个warp调度器每个调度器一个周期最多发射一条warp指令。理论峰值下这些发射槽应该几乎全部被矩阵指令或者配合矩阵计算的指令占据。但实际跑一个融合了大量标量计算的算子时发射槽会被地址计算、比较、跳转这些非计算指令大量挤占真正给到Tensor Core的HMMA指令占比自然就下来了。这种情况下单看TFLOPS完全无法反映真实性能。2. 指令发射单元的本质操作数检查、资源仲裁与调度策略2.1 发射前必须过的三重校验上一节提到发射前要检查三类情况这里展开说因为理解这些检查项后面读profile报告时才知道每个stall到底在等什么。第一重是操作数就绪检查主要处理的是数据依赖。最常见的是RAW依赖也就是读后写地址1的指令要读寄存器R1地址2的指令要写R1如果地址1那条还没执行完R1里的值就不能被读取。顺序发射的处理器一般用一张scoreboard表来记录每个寄存器正在被哪条指令占用发射时查一下就知道行不行乱序发射的处理器则是把等待中的指令放进保留站每个周期扫描一遍谁的操作数齐了就发射谁GPU的做法又不一样它靠warp级的切换来掩盖依赖——这个warp在等数据就换另一个warp来发射。所以GPU对依赖的容忍度更高代价是必须有足够多的warp来填坑。第二重是执行资源检查。执行单元是有限资源比如矩阵单元可能只有一套向量ALU有几条流水线每条流水线每个周期能接收几条指令是固定的。如果新指令的目标执行单元还在忙即使操作数就绪也得排队。还有一些更细的结构冲突比如某些芯片的浮点单元和整数单元共用写回端口或者特殊函数单元共享某个流水线阶段都会造成发射等待。第三重是端口冲突检查。寄存器堆的读端口是有限的比如物理上只有两个读端口那一个周期最多同时读两个寄存器操作数。如果发射宽度做得很大三条指令同时想读寄存器端口就成了瓶颈。访存端口也一样load/store单元能同时处理的事务数量有限超过就要stall。这一连串检查在硬件里是组合逻辑在软件里对应的是profile报告里那些stall原因。所以下次看到stall_wait要知道是依赖没解除看到stall_math_pipe_throttle是执行流水线本身拥塞看到stall_short_sb大概率是store缓存满了访存路径堵住了。2.2 顺序发射、乱序发射与协同发射AI芯片为何普遍选择简化围绕发射策略业界其实走过好几条路线。最原始的是顺序发射。指令按程序顺序一条一条检查、一条一条发遇到依赖就停下等前面指令完成再继续。好处是硬件极其简单坏处是性能波动大编译器很难填满所有停顿周期。今天基本只有低功耗嵌入式核心还在这么干。然后是乱序发射。以Tomasulo算法和保留站为代表等待发射的指令都挂在保留站里硬件持续扫描只要某条指令的操作数都到了执行单元也空着就把它提前发出去。这能让单线程的性能提升非常明显但代价是要额外做寄存器重命名、维护ROB重排序缓冲和物理寄存器堆面积和功耗都很高。CPU大核为了单线程性能愿意付出这个代价AI芯片一般付不起。还有一种思路是协同发射也就是VLIW风格。一条很宽的指令字里包含多个操作字段每个字段对应一类执行单元。发射动作本身极其简单指令到了按字段直接分发给对应的单元就行了因为编译器在生成这条指令时已经保证各个字段之间不冲突、不依赖。控制逻辑几乎为零。AI芯片特别吃这套因为它追求的是能效比和面积利用效率。代价也很明显——所有压力都推到编译器身上编译器弱一点性能就崩给你看。现在大多数AI加速芯片走的都是“简化硬件发射编译器强调度”的路线GPU则走了折中路线编译器生成相对规整的指令流硬件保留简单的调度器做warp级选择。理解了这些你就能明白为什么同一套模型在不同芯片上性能差距可以如此巨大——芯片的发射哲学不同对编译器的依赖程度完全不同。2.3 GPU warp scheduler 与 NPU VLIW 的对比我把两种主流发射架构放在一个表里对比看起来更直观。对比维度GPU warp schedulerNPU VLIW/多发射决策主体硬件调度器动态选warp编译器静态编排指令字段延迟隐藏靠大量warp切换靠指令级并行和DMA重叠发射宽度每个调度器通常1条/周期一条指令同时驱动多类单元分支处理硬件调度灵活支持动态分支编译器预判分支代价高代表产品NVIDIA SM内多个调度器TPU、多数ASIC NPUGPU的调度器更像一个灵活的值班交警现场看哪个warp能走就往哪个方向放NPU的VLIW则更像一个火车站时刻表车几点发、走哪条线都是提前排好的。火车站式的设计在列车准点编译器能把指令调好时效率极高一旦有晚点运行时有动态数据依赖或分支变化整个流水线就可能停滞而且没有硬件兜底。这也是为什么很多NPU跑固定shape的卷积、矩阵乘时性能惊艳一跑动态shape的Transformer就露馅。动态shape意味着每次输入的序列长度、batch大小都不一样编译器没法提前把每个循环的边界和双缓冲安排到最优指令流里就塞了一堆判断和分支发射槽利用率立刻掉下来。3. AI负载给发射单元的三重压力矩阵、访存与控制流混在一起3.1 GEMM的尺寸直接决定发射压力AI负载的主体是矩阵乘但矩阵乘和矩阵乘之间差别大了去了。一个明显的分界线是GEMM尺寸对发射开销的摊薄效应。训练阶段的大GEMM比如M、N、K都在4096以上每个矩阵乘指令对应的计算量非常大。拿HMMA指令来说一次可以处理一个16x16x16的矩阵乘分块里面包含几千次乘加操作。即使控制开销和访存开销存在摊到每条指令上发射端忙忙碌碌也好过那些算一次还用不了一两个周期的指令——因为一条指令就能喂饱MAC阵列很久。推理阶段的decode就完全是另一回事。大模型每个token生成时计算形态基本上从GEMM退化成了GEMV矩阵乘中的一个维度M变成1每条指令要算的数据量急剧减少与此同时负责搬运数据、计算地址、组织循环的标量指令比例大幅上升。这些指令一条条挤在发射槽里真正做矩阵计算的指令占比变得很低。A100跑大GEMM的Tensor Core利用率能到百分之七八十跑单token decode可能只有百分之十几差异就是这么来的。这也是我在管理算力平台时反复跟团队强调的一点看到GPU利用率低先别急着怪卡不行先看清楚负载形态是GEMM还是GEMV、batch是大是小。同样的卡跑连续批处理和跑在线推理呈现出来的利用率可以差三倍以上但卡本身一点问题都没有。3.2 访存指令和同步指令发射队列里的“无效占用”矩阵指令是计算的主力但AI kernel里还塞了大量访存和同步指令。这些指令不产生计算却在发射端占着宝贵的槽位有时还能把整条流水线堵死。load/store指令是最典型的。当一个kernel要频繁读写全局内存或共享内存时load/store单元的队列会很快占满。profiler里有个stall_short_sb意思是store buffer满了后续访存指令进不去依赖这些访存结果的计算指令更不可能发射。这种情况在瓶颈在内存带宽的算子中特别常见表现形式就是SM并不闲但Issue Active很低——大家不是不想干是干活的原料运不过来。barrier指令更狠。它要求同一块内的所有线程都执行到这个点才能继续只要组内最慢的那个线程还没到其他先到的线程就算操作数都就绪了也必须停在那里。发射端这时是完全空转的。同步点太多、同步粒度太粗都会让指令发射单元持续“罢工”。我处理过的一个经典问题某个NPU算子里编译后的指令流中DMA搬运和矩阵计算完全没有重叠每一块数据都要等DMA搬完、做一次barrier然后再启动计算算完再搬结果。在这整个过程中发射单元大部分时间在空转。后来在调度上改成双缓冲让下一块数据的DMA和当前块的计算重叠kernel耗时直接降了40%。你肉眼看不出指令流变化但profile一抓Issue Active从三十几提升到了七十多。3.3 稀疏计算和Tensor Core对指令级配合的依赖现在很多AI芯片标称算力动辄翻番靠的是稀疏计算支持常见的是2:4结构化稀疏也就是把权重里每四个连续元素强制压成两个非零配合索引元数据跳过一半乘加。算力翻倍的说法在MAC阵列层面是成立的但真正要吃到这波红利指令级配合有没有到位很关键。首先权重必须是稀疏格式且索引和数据要组织好才能被硬件识别其次每条稀疏指令要额外携带掩码信息说明哪些位置要计算、哪些跳过第三稀疏矩阵乘的访存模式不再规则数据搬运指令的发射节奏会比稠密更复杂。任何一个环节没跟上发射端就会被索引处理和数据重排拖住标称的2倍加速实际只有1.1倍甚至更差。Tensor Core的配合问题类似。一次HMMA指令确实能顶很多FFMA但它的前提是操作数已经通过ldmatrix等指令从共享内存搬到了寄存器。实际指令序列往往是ldmatrix、ldmatrix、HMMA、HMMA交替出现。如果发射端调度不好ldmatrix太慢HMMA就会因为操作数未就绪而等待或者HMMA发射太密矩阵单元忙不过来后面又堵住。要让Tensor Core持续满载发射端必须像节拍器一样稳定地踩准ldmatrix和HMMA之间的节奏。4. 发射逻辑不便宜面积、功耗与芯片厂商的取舍4.1 发射队列、比较器网络与寄存器堆端口聊完原理得回到芯片本身。很多人有一个错觉觉得指令发射单元只是一小片控制逻辑芯片面积和功耗的大头应该是那些MAC阵列和片上缓存。这个认知在AI芯片里有偏差在通用CPU里更是不成立。乱序发射核心的保留站和发射队列非常吃资源。每个发射队列项要保存操作码、源寄存器编号、目标寄存器编号、就绪状态位。更致命的是新指令进入队列时需要和队列里已有指令做寄存器依赖比较这个比较器网络是近似O(n^2)增长的。发射队列越深、发射宽度越宽比较器的面积和功耗增长越夸张。所以CPU大核为了做到6到8发射核心面积被撑得很大频率还压不下来。寄存器堆又是另一个无底洞。发射宽度每加宽寄存器堆的读写端口数量就要跟着加每个端口就是一堆位线和前置放大器布线成本极高。这也是为什么很多AI芯片的发射宽度并不宽但每一个发射槽都极其高效——它们宁愿把发射能力限定在“每周期发一条宽指令”也不愿意做成八路乱序发射。你去看物理设计报告会发现一个典型的现代高性能通用核心发射和重命名相关的逻辑可能占到核心面积的一到两成。AI芯片里这部分比例小一些但绝不是可以忽略的一块。厂商在设计时很清楚多省出一块面积就能多塞一组MAC阵列标称算力数字更好看市场宣传更有底气。4.2 脉动阵列的减法思路脉动阵列是另一个极端——它干脆把发射逻辑的复杂度降到最低。以TPU为代表MAC阵列里的每个处理单元直接从相邻单元拿数据不需要每个周期做一个大广播也不需要复杂的寄存器堆。控制逻辑和指令发射变得非常简单因为数据流是预先编排好的指令只需要告诉阵列“开始计算某一块”即可。这种“减法”思路的收益非常明显同样功耗预算下能塞进去的MAC数量更多能效比在规则大矩阵乘场景下做到极致。但代价同样明显——灵活性差。一旦碰到不规则计算脉动阵列的边界浪费就非常严重编译器得费很大力气去做分块、填充和映射。同一个公司后来推出的多核方案也会在脉动阵列边上加向量单元和标量单元就是为了让编译器有多类执行单元可以联合调度。对用户来说这意味着如果你要跑的工作负载是固定shape、大规模、规则矩阵乘脉动阵列风格的芯片性价比极高如果工作负载是动态shape、小算子为主、涉及大量控制流发射灵活性更重要脉动阵列反而帮不上忙。4.3 能效比引导下的发射宽度设计芯片厂商在规划下一代产品时发射宽度不是拍脑袋定的而是跟着能效比目标走。数据中心芯片希望每瓦算力越高越好那发射逻辑就得尽量精简把调度的重担交给编译器边缘推理芯片面对的模型五花八门泛化性要求高不能只靠编译器预排硬件调度能力就得保留一部分。我观察到的普遍趋势是通用GPU正在微调发射策略AI专用芯片则分成两派——一派把发射宽度和编译器配合做到极致一派在发射逻辑里塞进更多动态支持代价是能效轻微下降。选型时如果只看TOPS标称很容易选到前者跑你的模型时被动态shape打回原形。看芯片一定要看它的发射架构和软件栈成熟度同一颗芯片编译器排得好的kernel和排不好的kernel吞吐能差几倍。5. 大模型推理的真实profile发射端空转才是算力利用率低的元凶5.1 用Nsight Compute看一次推理任务从No Eligible到warp stall讲完理论落地到工具。我平时在NVIDIA卡上定位发射瓶颈基本走三步。先跑nsys做一次整体profile找出耗时最靠前的kernel。不要一上来就ncu全量抓数据量巨大且浪费时间。接着针对Top kernel执行ncu --set full --section SchedulerStats --section WarpStateStats --launch-count 3 ./your_application或者二次采样时用ncu --kernel-name regex --launch-skip 5 --launch-count 3 --section SchedulerStats --section WarpStateStats ./your_application重点看SchedulerStats里的Issue Active和No Eligible。Issue Active是调度器真正发射指令的周期占比No Eligible是调度器空转、找不到可发射warp的周期占比。decode类kernel的profile里No Eligible经常能占到三成甚至五成也就是说调度器有一半时间在干瞪眼。再结合WarpStateStats里的stall原因分布基本就能定位如果是stall_wait占比高说明依赖写回太慢运算单元和访存路径要多检查如果是stall_short_sb高数据搬回来之前store buffer先满了重点是访存调度如果是stall_drain高主要出现在kernel收尾所有warp都在等最后的指令执行完写回可能和并行度分配不均衡有关。5.2 token吞吐、KV Cache访存与发射效率的三角关系大模型推理时token吞吐这个概念大家都不陌生。但真正评估一个芯片能跑多少token每秒不能只算FLOPs要同时看三个约束算力、访存带宽、指令发射效率。当batch很小比如在线推理常遇到的batch1计算量极低但每个token都要读一遍KV Cache里所有历史token对应的键值向量访存压力巨大。此时计算单元经常空转发射端则在反复等待数据回来。反过来当batch拉到64或128同一份权重被更多token共享GEMM尺寸变大算力逐渐变成主要瓶颈发射效率和访存压力反而下降。这就是为什么连续批处理里加大batch能显著提升卡利用率。这里我想分享一个真实估算翻车的案例。有次我给团队估算7B模型的推理吞吐按FLOPs算理论上限应该是每秒200多个token结果实测只有30多。事后复盘发现decode阶段的访存需求我按KV Cache访问量算对了但没把指令发射效率算进去——大量小矩阵乘和访存指令把发射槽占满真正能发射矩阵指令的周期少得可怜。从那以后我评估token算力需求时都看两组数字一组按FLOPs推的理论吞吐一组按访存带宽和指令条数推的保守吞吐两者之间取低值做参考。5.3 同一个算子不同尺寸下的发射级表现还有一个容易被忽略的场景同一个算子因为输入尺寸不同吃掉的瓶颈资源完全不同。最典型的就是注意力里的QK^T。当序列长度很短时QK^T是一个小矩阵乘指令发射和kernel launch的开销占比很高此时无论你怎么优化MAC阵列性能都上不去瓶颈在发射端和访存延迟。当序列长度很长时QK^T变成大GEMM矩阵指令占比上来了算力才真正成为瓶颈。同样是“矩阵乘”三个字发射端的压力天差地别。遇到这种场景我建议用固定模型、改变输入序列长度的方法做scaling测试每次只改一个变量跑ncu看Issue Active和No Eligible随长度的变化曲线。这个曲线能直接告诉你候选芯片在哪个长度区间能吃满算力在哪个区间只能干瞪眼。做算力平台选型时把这套scaling测试脚本沉淀下来比看任何官方benchmark都有用。6. 从单卡到集群算力管理员如何观测发射瓶颈并做选型6.1 定位发射瓶颈的操作路径前面提到的工作流完整串起来是这样的先用nvidia-smi dmon这类工具看整卡宏观指标确认没有别的进程抢占SM利用率、显存利用率、功耗是不是正常。然后跑一遍PyTorch Profiler或者TensorBoard Profiler拿到kernel耗时排名挑出top 1到3个kernel。最后只对这几个关键kernel做ncu深度分析看Issue Active、No Eligible和stall分布。这个顺序千万别反。直接跳过前面两层去跑ncu很常见结果就是抓到一些无关紧要的kernel上浪费大半天时间。正确的排错逻辑是先确认“整体上到底哪个阶段最慢”再深挖“这个阶段为什么慢”。6.2 多台异构算力服务器的统一观测指标管理多台算力服务器尤其是混部了不同品牌、不同架构的GPU或AI加速卡时统一指标口径很重要。我自己用DCGM exporter加Prometheus加Grafana这套栈把每台机器的关键指标汇总成一个面板。推荐的指标优先级是SM Active、Tensor Pipe Active、DRAM Active再加显存时钟和功耗。为什么不是只看SM Active因为SM Active只表示SM上有指令在活跃执行但执行的可能是一堆地址计算和非计算指令不代表真正在运算。Tensor Pipe Active能更准确地反映矩阵单元的真实占用率DRAM Active则反映访存压力。这三个放在一起基本能判断一个kernel是算力受限、访存受限还是发射受限。注意SM Active高和算力利用率高是两回事。曾经有个跨部门协作的项目对方拿着一份SM Active 95%的报告说芯片已经跑满了结果我们打开Perfetto一看一大半周期都耗在自旋等待上。SM Active只是“SM没闲着”不等于“SM都在干正事”。另外异构卡混部时我建议按“指令发射表现”而不是按品牌型号来分组。同一个大模型任务放到A型卡上Issue Active能到70%放到B型卡上只有30%即使两者标称TOPS一样跑起来的时间和耗电也差很远。统一管理平台里把每个kernel的发射效率统计好可以作为调度器的软亲和依据把任务尽量调度到发射效率高的卡上。6.3 算力选型与平台管理经验清单最后整理一份我积累下来的经验清单算是给新入行的人抄作业不要拿纸面TOPS或TFLOPS排行直接决定购买。同一组数字背后可能是完全不同的发射架构、软件栈和编译器成熟度。选型测试要覆盖三类典型负载大GEMM、小GEMV、注意力多尺寸。分别看算力专用单元的利用率和Issue Active。关注软件栈对动态shape的支持程度。动态shape下编译器能不能快速生成高效的指令流直接决定线上推理的尾延迟和平均吞吐。如果是租算力平台务必确认能否拿到profiler权限。拿不到Nsight Compute或类似工具的权限就很难判断性能瓶颈出了问题只能黑盒猜测。内部建立算力基准集用你自己真实模型的代表性算子定期跑观察Issue Active和No Eligible的波动。基线一旦建立后续芯片升级、驱动更新、框架版本变化都能快速做对比。我在实际管平台时感受最深的一点是指令发射单元的瓶颈往往不在芯片宣传页上也不在TOPS排行里但它确确实实决定了算力天花板的天花板——就算你把显存加满、网络改好、调度器调优到极限发射端喂不进去指令一切都白搭。理解这个十字路口是每一个做算力运营和AI Infra的人绕不开的功课。 指令发射单元：AI芯片真实算力吞吐的隐形瓶颈
很多做AI Infra和算力运营的朋友应该都碰到过这种怪事租来的GPU明明标称几百T算力跑大模型推理时nvidia-smi里看到的利用率却只有百分之二三十跑矩阵乘benchmark又能跑出非常漂亮的数字。差距到底出在哪我刚接手算力平台那会儿也被这个问题折磨过后来把底层性能分析彻底摸了一遍才发现大部分人忽略了一个微架构层面的关键角色——指令发射单元。它是AI芯片的“十字路口”每条指令从取指译码后都必须经过这里才能进入真正干活的执行单元。这个路口的调度能力、等待队列和拥堵状况决定了纸面上的算力指标能转化出多少真实吞吐也就是我们常说的算力天花板。这篇文章适合谁看一方面是在做芯片选型、算力平台管理、大模型推理优化的工程师另一方面是纯粹想搞明白“为什么我的卡利用率上不去”的算法同学。我会从微架构原理讲起穿插我实际在GPU和各类AI加速卡上做profile的经验最后落到多卡集群管理层面尽量让不同背景的人都能照着用。1. 标称算力只是理论值指令发射效率才是真实吞吐的阀门1.1 一条指令从取指到执行要过几道门我们先从最基础的过程看起。一段神经网络模型经过编译器最终会变成一连串机器指令。这些指令要先被取指单元从一级指令缓存里一条条取出来送到译码器。译码器做的事情是把操作码、寄存器地址、立即数拆开搞清楚这条指令到底想干什么、要读哪些寄存器、要写到哪去。到这一步为止很多人的理解就结束了总觉得编译完的指令会顺理成章地执行。实际上不是。译码完成之后的指令必须进入指令发射单元由它来判定这条指令能不能立刻往下走。要检查的东西至少有三项操作数是不是已经就绪——比如上一条指令刚说要写某个寄存器结果还没写回来这条指令就去读那肯定不行执行单元是不是有空闲——矩阵单元还在跑上一条大矩阵乘新的矩阵指令就只能等着寄存器和访存端口有没有冲突——有的芯片寄存器堆读写端口有限两条指令同时要读同一组端口总得有个先来后到。全部通过之后指令才被真正“发射”到执行单元上。这个过程用交通来类比特别贴切。执行单元是城市里的各条车道指令发射单元就是进车道之前的十字路口。路口每个周期能放行多少辆车决定了后面车道修得再宽、跑车性能再好都没用。很多时候我们嫌某个AI芯片“跑不出标称算力”本质不是后面的MAC阵列不够快而是这个路口在过去车。1.2 IPC、发射宽度和发射槽利用率衡量路口通畅程度的三把尺子要量化这个路口的通行能力有几个基本指标得认识。第一个是IPC平均每个时钟周期完成或写回的指令数。老玩家看CPU评测时总能看到这个数字它反映的是整个从取指到执行的回流是否顺畅。IPC过低说明指令流里存在大量等待发射端和计算端没有咬合好。第二个是发射宽度也就是每个时钟周期最多能发射几条指令。这里的“条”在不同架构里含义不同。在GPU里指一个调度器每周期可以发射一条warp指令一条warp指令对应32条线程的同一操作在VLIW风格的NPU里指一条超长指令字里有几个作用域可以同时触发比如同时控制矩阵单元、向量单元和DMA引擎。第三个也是我平时最关注的发射槽利用率。它等于实际发出的指令数除以理论可发射的槽位数。很多kernel跑下来这个值可能只有40%到50%意味着每两个周期里有一个周期的发射机会是彻彻底底空着的。GPU的profiler里有个专门的统计项叫Issue Active说的就是这个比例还有一个对应的No Eligible表示调度器想发射但找了一圈发现所有warp都不满足发射条件只能空转。我自己的经验是这三把尺子得配合着看。SM的占用率很高、但Issue Active很低说明执行单元虽然都在线发射端根本没有把计算指令喂进来大概率是访存或者依赖把warp卡住了如果SM占用率也很低那就可能是kernel启动不足比如batch太小整个SM里的线程块没排满。1.3 为什么标称TFLOPS只是理论上限标称算力这个概念本身有个隐含假设每个时钟周期芯片上所有的矩阵单元都在满负荷跑矩阵运算。这个假设要成立必须满足两个条件。第一矩阵指令是连续不断送到MAC阵列的中间不能有断流第二每条矩阵指令要读的数据已经全部准备好放到共享内存或者寄存器里了。第一个条件测的就是指令发射单元。只要发射断流不管标称TFLOPS是300还是1000MAC阵列都只能原地等待。第二个条件其实也间接和发射有关——如果数据搬运指令没有提前发射出去后面的计算指令即便已经轮到发射了也会因为操作数没就绪而发不出去。举个例子A100的SM里有4个warp调度器每个调度器一个周期最多发射一条warp指令。理论峰值下这些发射槽应该几乎全部被矩阵指令或者配合矩阵计算的指令占据。但实际跑一个融合了大量标量计算的算子时发射槽会被地址计算、比较、跳转这些非计算指令大量挤占真正给到Tensor Core的HMMA指令占比自然就下来了。这种情况下单看TFLOPS完全无法反映真实性能。2. 指令发射单元的本质操作数检查、资源仲裁与调度策略2.1 发射前必须过的三重校验上一节提到发射前要检查三类情况这里展开说因为理解这些检查项后面读profile报告时才知道每个stall到底在等什么。第一重是操作数就绪检查主要处理的是数据依赖。最常见的是RAW依赖也就是读后写地址1的指令要读寄存器R1地址2的指令要写R1如果地址1那条还没执行完R1里的值就不能被读取。顺序发射的处理器一般用一张scoreboard表来记录每个寄存器正在被哪条指令占用发射时查一下就知道行不行乱序发射的处理器则是把等待中的指令放进保留站每个周期扫描一遍谁的操作数齐了就发射谁GPU的做法又不一样它靠warp级的切换来掩盖依赖——这个warp在等数据就换另一个warp来发射。所以GPU对依赖的容忍度更高代价是必须有足够多的warp来填坑。第二重是执行资源检查。执行单元是有限资源比如矩阵单元可能只有一套向量ALU有几条流水线每条流水线每个周期能接收几条指令是固定的。如果新指令的目标执行单元还在忙即使操作数就绪也得排队。还有一些更细的结构冲突比如某些芯片的浮点单元和整数单元共用写回端口或者特殊函数单元共享某个流水线阶段都会造成发射等待。第三重是端口冲突检查。寄存器堆的读端口是有限的比如物理上只有两个读端口那一个周期最多同时读两个寄存器操作数。如果发射宽度做得很大三条指令同时想读寄存器端口就成了瓶颈。访存端口也一样load/store单元能同时处理的事务数量有限超过就要stall。这一连串检查在硬件里是组合逻辑在软件里对应的是profile报告里那些stall原因。所以下次看到stall_wait要知道是依赖没解除看到stall_math_pipe_throttle是执行流水线本身拥塞看到stall_short_sb大概率是store缓存满了访存路径堵住了。2.2 顺序发射、乱序发射与协同发射AI芯片为何普遍选择简化围绕发射策略业界其实走过好几条路线。最原始的是顺序发射。指令按程序顺序一条一条检查、一条一条发遇到依赖就停下等前面指令完成再继续。好处是硬件极其简单坏处是性能波动大编译器很难填满所有停顿周期。今天基本只有低功耗嵌入式核心还在这么干。然后是乱序发射。以Tomasulo算法和保留站为代表等待发射的指令都挂在保留站里硬件持续扫描只要某条指令的操作数都到了执行单元也空着就把它提前发出去。这能让单线程的性能提升非常明显但代价是要额外做寄存器重命名、维护ROB重排序缓冲和物理寄存器堆面积和功耗都很高。CPU大核为了单线程性能愿意付出这个代价AI芯片一般付不起。还有一种思路是协同发射也就是VLIW风格。一条很宽的指令字里包含多个操作字段每个字段对应一类执行单元。发射动作本身极其简单指令到了按字段直接分发给对应的单元就行了因为编译器在生成这条指令时已经保证各个字段之间不冲突、不依赖。控制逻辑几乎为零。AI芯片特别吃这套因为它追求的是能效比和面积利用效率。代价也很明显——所有压力都推到编译器身上编译器弱一点性能就崩给你看。现在大多数AI加速芯片走的都是“简化硬件发射编译器强调度”的路线GPU则走了折中路线编译器生成相对规整的指令流硬件保留简单的调度器做warp级选择。理解了这些你就能明白为什么同一套模型在不同芯片上性能差距可以如此巨大——芯片的发射哲学不同对编译器的依赖程度完全不同。2.3 GPU warp scheduler 与 NPU VLIW 的对比我把两种主流发射架构放在一个表里对比看起来更直观。对比维度GPU warp schedulerNPU VLIW/多发射决策主体硬件调度器动态选warp编译器静态编排指令字段延迟隐藏靠大量warp切换靠指令级并行和DMA重叠发射宽度每个调度器通常1条/周期一条指令同时驱动多类单元分支处理硬件调度灵活支持动态分支编译器预判分支代价高代表产品NVIDIA SM内多个调度器TPU、多数ASIC NPUGPU的调度器更像一个灵活的值班交警现场看哪个warp能走就往哪个方向放NPU的VLIW则更像一个火车站时刻表车几点发、走哪条线都是提前排好的。火车站式的设计在列车准点编译器能把指令调好时效率极高一旦有晚点运行时有动态数据依赖或分支变化整个流水线就可能停滞而且没有硬件兜底。这也是为什么很多NPU跑固定shape的卷积、矩阵乘时性能惊艳一跑动态shape的Transformer就露馅。动态shape意味着每次输入的序列长度、batch大小都不一样编译器没法提前把每个循环的边界和双缓冲安排到最优指令流里就塞了一堆判断和分支发射槽利用率立刻掉下来。3. AI负载给发射单元的三重压力矩阵、访存与控制流混在一起3.1 GEMM的尺寸直接决定发射压力AI负载的主体是矩阵乘但矩阵乘和矩阵乘之间差别大了去了。一个明显的分界线是GEMM尺寸对发射开销的摊薄效应。训练阶段的大GEMM比如M、N、K都在4096以上每个矩阵乘指令对应的计算量非常大。拿HMMA指令来说一次可以处理一个16x16x16的矩阵乘分块里面包含几千次乘加操作。即使控制开销和访存开销存在摊到每条指令上发射端忙忙碌碌也好过那些算一次还用不了一两个周期的指令——因为一条指令就能喂饱MAC阵列很久。推理阶段的decode就完全是另一回事。大模型每个token生成时计算形态基本上从GEMM退化成了GEMV矩阵乘中的一个维度M变成1每条指令要算的数据量急剧减少与此同时负责搬运数据、计算地址、组织循环的标量指令比例大幅上升。这些指令一条条挤在发射槽里真正做矩阵计算的指令占比变得很低。A100跑大GEMM的Tensor Core利用率能到百分之七八十跑单token decode可能只有百分之十几差异就是这么来的。这也是我在管理算力平台时反复跟团队强调的一点看到GPU利用率低先别急着怪卡不行先看清楚负载形态是GEMM还是GEMV、batch是大是小。同样的卡跑连续批处理和跑在线推理呈现出来的利用率可以差三倍以上但卡本身一点问题都没有。3.2 访存指令和同步指令发射队列里的“无效占用”矩阵指令是计算的主力但AI kernel里还塞了大量访存和同步指令。这些指令不产生计算却在发射端占着宝贵的槽位有时还能把整条流水线堵死。load/store指令是最典型的。当一个kernel要频繁读写全局内存或共享内存时load/store单元的队列会很快占满。profiler里有个stall_short_sb意思是store buffer满了后续访存指令进不去依赖这些访存结果的计算指令更不可能发射。这种情况在瓶颈在内存带宽的算子中特别常见表现形式就是SM并不闲但Issue Active很低——大家不是不想干是干活的原料运不过来。barrier指令更狠。它要求同一块内的所有线程都执行到这个点才能继续只要组内最慢的那个线程还没到其他先到的线程就算操作数都就绪了也必须停在那里。发射端这时是完全空转的。同步点太多、同步粒度太粗都会让指令发射单元持续“罢工”。我处理过的一个经典问题某个NPU算子里编译后的指令流中DMA搬运和矩阵计算完全没有重叠每一块数据都要等DMA搬完、做一次barrier然后再启动计算算完再搬结果。在这整个过程中发射单元大部分时间在空转。后来在调度上改成双缓冲让下一块数据的DMA和当前块的计算重叠kernel耗时直接降了40%。你肉眼看不出指令流变化但profile一抓Issue Active从三十几提升到了七十多。3.3 稀疏计算和Tensor Core对指令级配合的依赖现在很多AI芯片标称算力动辄翻番靠的是稀疏计算支持常见的是2:4结构化稀疏也就是把权重里每四个连续元素强制压成两个非零配合索引元数据跳过一半乘加。算力翻倍的说法在MAC阵列层面是成立的但真正要吃到这波红利指令级配合有没有到位很关键。首先权重必须是稀疏格式且索引和数据要组织好才能被硬件识别其次每条稀疏指令要额外携带掩码信息说明哪些位置要计算、哪些跳过第三稀疏矩阵乘的访存模式不再规则数据搬运指令的发射节奏会比稠密更复杂。任何一个环节没跟上发射端就会被索引处理和数据重排拖住标称的2倍加速实际只有1.1倍甚至更差。Tensor Core的配合问题类似。一次HMMA指令确实能顶很多FFMA但它的前提是操作数已经通过ldmatrix等指令从共享内存搬到了寄存器。实际指令序列往往是ldmatrix、ldmatrix、HMMA、HMMA交替出现。如果发射端调度不好ldmatrix太慢HMMA就会因为操作数未就绪而等待或者HMMA发射太密矩阵单元忙不过来后面又堵住。要让Tensor Core持续满载发射端必须像节拍器一样稳定地踩准ldmatrix和HMMA之间的节奏。4. 发射逻辑不便宜面积、功耗与芯片厂商的取舍4.1 发射队列、比较器网络与寄存器堆端口聊完原理得回到芯片本身。很多人有一个错觉觉得指令发射单元只是一小片控制逻辑芯片面积和功耗的大头应该是那些MAC阵列和片上缓存。这个认知在AI芯片里有偏差在通用CPU里更是不成立。乱序发射核心的保留站和发射队列非常吃资源。每个发射队列项要保存操作码、源寄存器编号、目标寄存器编号、就绪状态位。更致命的是新指令进入队列时需要和队列里已有指令做寄存器依赖比较这个比较器网络是近似O(n^2)增长的。发射队列越深、发射宽度越宽比较器的面积和功耗增长越夸张。所以CPU大核为了做到6到8发射核心面积被撑得很大频率还压不下来。寄存器堆又是另一个无底洞。发射宽度每加宽寄存器堆的读写端口数量就要跟着加每个端口就是一堆位线和前置放大器布线成本极高。这也是为什么很多AI芯片的发射宽度并不宽但每一个发射槽都极其高效——它们宁愿把发射能力限定在“每周期发一条宽指令”也不愿意做成八路乱序发射。你去看物理设计报告会发现一个典型的现代高性能通用核心发射和重命名相关的逻辑可能占到核心面积的一到两成。AI芯片里这部分比例小一些但绝不是可以忽略的一块。厂商在设计时很清楚多省出一块面积就能多塞一组MAC阵列标称算力数字更好看市场宣传更有底气。4.2 脉动阵列的减法思路脉动阵列是另一个极端——它干脆把发射逻辑的复杂度降到最低。以TPU为代表MAC阵列里的每个处理单元直接从相邻单元拿数据不需要每个周期做一个大广播也不需要复杂的寄存器堆。控制逻辑和指令发射变得非常简单因为数据流是预先编排好的指令只需要告诉阵列“开始计算某一块”即可。这种“减法”思路的收益非常明显同样功耗预算下能塞进去的MAC数量更多能效比在规则大矩阵乘场景下做到极致。但代价同样明显——灵活性差。一旦碰到不规则计算脉动阵列的边界浪费就非常严重编译器得费很大力气去做分块、填充和映射。同一个公司后来推出的多核方案也会在脉动阵列边上加向量单元和标量单元就是为了让编译器有多类执行单元可以联合调度。对用户来说这意味着如果你要跑的工作负载是固定shape、大规模、规则矩阵乘脉动阵列风格的芯片性价比极高如果工作负载是动态shape、小算子为主、涉及大量控制流发射灵活性更重要脉动阵列反而帮不上忙。4.3 能效比引导下的发射宽度设计芯片厂商在规划下一代产品时发射宽度不是拍脑袋定的而是跟着能效比目标走。数据中心芯片希望每瓦算力越高越好那发射逻辑就得尽量精简把调度的重担交给编译器边缘推理芯片面对的模型五花八门泛化性要求高不能只靠编译器预排硬件调度能力就得保留一部分。我观察到的普遍趋势是通用GPU正在微调发射策略AI专用芯片则分成两派——一派把发射宽度和编译器配合做到极致一派在发射逻辑里塞进更多动态支持代价是能效轻微下降。选型时如果只看TOPS标称很容易选到前者跑你的模型时被动态shape打回原形。看芯片一定要看它的发射架构和软件栈成熟度同一颗芯片编译器排得好的kernel和排不好的kernel吞吐能差几倍。5. 大模型推理的真实profile发射端空转才是算力利用率低的元凶5.1 用Nsight Compute看一次推理任务从No Eligible到warp stall讲完理论落地到工具。我平时在NVIDIA卡上定位发射瓶颈基本走三步。先跑nsys做一次整体profile找出耗时最靠前的kernel。不要一上来就ncu全量抓数据量巨大且浪费时间。接着针对Top kernel执行ncu --set full --section SchedulerStats --section WarpStateStats --launch-count 3 ./your_application或者二次采样时用ncu --kernel-name regex --launch-skip 5 --launch-count 3 --section SchedulerStats --section WarpStateStats ./your_application重点看SchedulerStats里的Issue Active和No Eligible。Issue Active是调度器真正发射指令的周期占比No Eligible是调度器空转、找不到可发射warp的周期占比。decode类kernel的profile里No Eligible经常能占到三成甚至五成也就是说调度器有一半时间在干瞪眼。再结合WarpStateStats里的stall原因分布基本就能定位如果是stall_wait占比高说明依赖写回太慢运算单元和访存路径要多检查如果是stall_short_sb高数据搬回来之前store buffer先满了重点是访存调度如果是stall_drain高主要出现在kernel收尾所有warp都在等最后的指令执行完写回可能和并行度分配不均衡有关。5.2 token吞吐、KV Cache访存与发射效率的三角关系大模型推理时token吞吐这个概念大家都不陌生。但真正评估一个芯片能跑多少token每秒不能只算FLOPs要同时看三个约束算力、访存带宽、指令发射效率。当batch很小比如在线推理常遇到的batch1计算量极低但每个token都要读一遍KV Cache里所有历史token对应的键值向量访存压力巨大。此时计算单元经常空转发射端则在反复等待数据回来。反过来当batch拉到64或128同一份权重被更多token共享GEMM尺寸变大算力逐渐变成主要瓶颈发射效率和访存压力反而下降。这就是为什么连续批处理里加大batch能显著提升卡利用率。这里我想分享一个真实估算翻车的案例。有次我给团队估算7B模型的推理吞吐按FLOPs算理论上限应该是每秒200多个token结果实测只有30多。事后复盘发现decode阶段的访存需求我按KV Cache访问量算对了但没把指令发射效率算进去——大量小矩阵乘和访存指令把发射槽占满真正能发射矩阵指令的周期少得可怜。从那以后我评估token算力需求时都看两组数字一组按FLOPs推的理论吞吐一组按访存带宽和指令条数推的保守吞吐两者之间取低值做参考。5.3 同一个算子不同尺寸下的发射级表现还有一个容易被忽略的场景同一个算子因为输入尺寸不同吃掉的瓶颈资源完全不同。最典型的就是注意力里的QK^T。当序列长度很短时QK^T是一个小矩阵乘指令发射和kernel launch的开销占比很高此时无论你怎么优化MAC阵列性能都上不去瓶颈在发射端和访存延迟。当序列长度很长时QK^T变成大GEMM矩阵指令占比上来了算力才真正成为瓶颈。同样是“矩阵乘”三个字发射端的压力天差地别。遇到这种场景我建议用固定模型、改变输入序列长度的方法做scaling测试每次只改一个变量跑ncu看Issue Active和No Eligible随长度的变化曲线。这个曲线能直接告诉你候选芯片在哪个长度区间能吃满算力在哪个区间只能干瞪眼。做算力平台选型时把这套scaling测试脚本沉淀下来比看任何官方benchmark都有用。6. 从单卡到集群算力管理员如何观测发射瓶颈并做选型6.1 定位发射瓶颈的操作路径前面提到的工作流完整串起来是这样的先用nvidia-smi dmon这类工具看整卡宏观指标确认没有别的进程抢占SM利用率、显存利用率、功耗是不是正常。然后跑一遍PyTorch Profiler或者TensorBoard Profiler拿到kernel耗时排名挑出top 1到3个kernel。最后只对这几个关键kernel做ncu深度分析看Issue Active、No Eligible和stall分布。这个顺序千万别反。直接跳过前面两层去跑ncu很常见结果就是抓到一些无关紧要的kernel上浪费大半天时间。正确的排错逻辑是先确认“整体上到底哪个阶段最慢”再深挖“这个阶段为什么慢”。6.2 多台异构算力服务器的统一观测指标管理多台算力服务器尤其是混部了不同品牌、不同架构的GPU或AI加速卡时统一指标口径很重要。我自己用DCGM exporter加Prometheus加Grafana这套栈把每台机器的关键指标汇总成一个面板。推荐的指标优先级是SM Active、Tensor Pipe Active、DRAM Active再加显存时钟和功耗。为什么不是只看SM Active因为SM Active只表示SM上有指令在活跃执行但执行的可能是一堆地址计算和非计算指令不代表真正在运算。Tensor Pipe Active能更准确地反映矩阵单元的真实占用率DRAM Active则反映访存压力。这三个放在一起基本能判断一个kernel是算力受限、访存受限还是发射受限。注意SM Active高和算力利用率高是两回事。曾经有个跨部门协作的项目对方拿着一份SM Active 95%的报告说芯片已经跑满了结果我们打开Perfetto一看一大半周期都耗在自旋等待上。SM Active只是“SM没闲着”不等于“SM都在干正事”。另外异构卡混部时我建议按“指令发射表现”而不是按品牌型号来分组。同一个大模型任务放到A型卡上Issue Active能到70%放到B型卡上只有30%即使两者标称TOPS一样跑起来的时间和耗电也差很远。统一管理平台里把每个kernel的发射效率统计好可以作为调度器的软亲和依据把任务尽量调度到发射效率高的卡上。6.3 算力选型与平台管理经验清单最后整理一份我积累下来的经验清单算是给新入行的人抄作业不要拿纸面TOPS或TFLOPS排行直接决定购买。同一组数字背后可能是完全不同的发射架构、软件栈和编译器成熟度。选型测试要覆盖三类典型负载大GEMM、小GEMV、注意力多尺寸。分别看算力专用单元的利用率和Issue Active。关注软件栈对动态shape的支持程度。动态shape下编译器能不能快速生成高效的指令流直接决定线上推理的尾延迟和平均吞吐。如果是租算力平台务必确认能否拿到profiler权限。拿不到Nsight Compute或类似工具的权限就很难判断性能瓶颈出了问题只能黑盒猜测。内部建立算力基准集用你自己真实模型的代表性算子定期跑观察Issue Active和No Eligible的波动。基线一旦建立后续芯片升级、驱动更新、框架版本变化都能快速做对比。我在实际管平台时感受最深的一点是指令发射单元的瓶颈往往不在芯片宣传页上也不在TOPS排行里但它确确实实决定了算力天花板的天花板——就算你把显存加满、网络改好、调度器调优到极限发射端喂不进去指令一切都白搭。理解这个十字路口是每一个做算力运营和AI Infra的人绕不开的功课。