手里同时压着GPU服务器和NPU开发板的人大概率都经历过同一种分裂感同一段模型代码在GPU上跑得稳稳当当搬到NPU上不是算子不支持就是精度掉点再不然就是速度根本没达到标称算力。GPU、NPU、指令集、稀疏化、存算一体这五个词看着像是五个独立话题实际上它们是一条线上的五个环节——从最底层一条指令怎么被发射出去到数据怎么在存储单元和计算单元之间搬来搬去再到怎么用稀疏和存内计算绕开带宽这堵墙。这条线不走通选型就会变成看参数表拍脑袋。这篇东西写给三类人正在做推理加速、纠结该上GPU还是NPU的工程同学写过CUDA kernel、想搞明白NPU那套数据流到底差在哪的底层同学以及做模型部署、被量化精度和带宽瓶颈反复折磨的算法同学。我会尽量少讲结论、多讲推导过程把每个设计选择背后的账算给你看包括我自己在实测里踩过的那些坑。1. 先想清楚GPU和NPU解决的根本不是同一个问题1.1 一次模型迁移失败引出的追问我印象最深的一次是把一个带动态shape的检测模型往NPU上搬。GPU上这套东西跑得挺自在因为CUDA的编程模型本质上是给你一堆线程随便你写控制流靠分支掩码处理访存靠多级缓存兜底编译器再差也能跑起来最多是慢。到了NPU上问题变成另一副面孔编译器在构图阶段就直接报错因为它的指令系统里压根没有按运行时数据决定循环次数这种表达能力。这不是NPU做得差而是它从设计第一天起就没打算支持这种灵活性。这件事把双方的分野讲透了。GPU的出发点是我不知道你要算什么所以我把尽可能通用的能力塞进去NPU的出发点是我知道你八成在算卷积和矩阵乘那我就把这两件事做到极致其余的我选择放弃。前者用灵活性换确定性后者用确定性换效率。所有后续的差异——指令集形态、数据通路宽度、缓存层级、对稀疏的支持方式——都是这个初始选择的必然结果。理解了这一点你就不会再问NPU为什么不能跑我的自定义算子这种问题了。正确的问法是我这个算子能不能被拆解成NPU已有的原语组合如果拆不了那就得回到GPU或者改模型。1.2 通用性与专用性的取舍账拿一个具体的数字来感受。GPU上一个FP32 ALU做一次乘加理论上要占满流水线而NPU里一个INT8 MAC阵列同样面积上能塞进去的乘加单元数量往往高出一到两个数量级。原因不神秘GPU的每个计算单元后面拖着寄存器堆、指令缓存、调度逻辑、分支处理、地址生成这一整套通用机制这些开销在通用场景下是必要的在纯矩阵乘场景下就全是浪费。NPU把这些全部砍掉只留下读数据、乘加、写数据这条最短路径面积自然就省出来了。但省出来的面积有代价。GPU可以在同一个芯片上跑图形、跑渲染、跑科学计算、跑训练和推理靠的是那一整套通用机制NPU只能干它被设计出来的那件事设计时没考虑到的算子要么退化成低效的通用实现要么直接不支持。行业里那句老话——GPU是瑞士军刀NPU是专用扳手——说的就是这个。实际选型的时候如果你的模型结构稳定、算子集合固定、以推理为主NPU的能效比优势是真实的如果你的模型还在快速迭代、算子每周都在变、还要兼顾训练那GPU的通用性反而省下了大量的迁移成本。提醒一句能效比TOPS/W这个指标特别容易被拿来当卖点但它是峰值条件下的数字。实际部署里能不能逼近峰值取决于你的模型能不能在数据复用上喂饱那块阵列这个后面会详细算。2. 指令集层面SIMT与数据流架构的分野2.1 GPU的SIMT是怎么把一条指令铺到32个线程GPU的核心执行模型是SIMTSingle Instruction, Multiple Threads。拿一段最普通的向量加法举例你写的是标量代码编译器生成的却是一条指令驱动一整个warp通常是32个线程去执行。每个线程有自己的寄存器和程序计数器名义上的私有状态但硬件层面这些状态是被打包进寄存器堆、按warp为单位统一调度的。看一段真实的PTX会直观很多// 每个线程根据自己的 tid 取出一个元素 ld.global.nc.v4.f32 {%f1,%f2,%f3,%f4}, [%rd1]; // 向量化乘加四个分量一次性算完 fma.rn.f32 %f5, %f1, %f2, %f5; fma.rn.f32 %f6, %f3, %f4, %f6; st.global.v4.f32 [%rd2], {%f5,%f6,%f7,%f8};关键在于这三条指令是整个warp一起走的。32个线程各自拿着自己的地址、各自算自己的结果但指令只有一份。这就是SIMT省下来的东西——指令获取和解码的开销被32个线程摊薄了。那分支怎么办如果warp里的线程走进不同的分支硬件不会给每个线程单独发指令而是把两条路径串行执行两遍用执行掩码把不该走这条路的线程屏蔽掉。极端情况下一个32线程的warp里出现32个不同分支执行效率直接掉到1/32。这不是bug是SIMT模型与生俱来的代价。所以GPU优化里有一条铁律尽量让同一个warp里的线程访问连续地址、走同一条控制流。我把这条规则写在团队文档第一行因为它是所有后续优化的前提。2.2 warp、CTA、SM这三级组织为什么这样设计很多人第一次看到CUDA的线程层级会觉得冗余线程、warp、blockCTA、grid为什么搞这么多层。其实每一层都对应一个明确的硬件资源分配边界。warp是调度最小单位一个warp占一份指令发射槽CTA是资源分配单位一个CTA会被整体放到一个SM上占用的寄存器、共享内存、warp槽位在它存续期间不会被别的CTA抢走SM是物理计算单元内部有多个调度器和四类执行单元FP32、INT32、特殊功能、加载存储。这套组织的设计意图是让程序员能用局部性换性能。同一个CTA里的warp可以通过共享内存交换数据不需要走全局显存共享内存的延迟大概比L2低一个数量级比HBM低两个数量级。矩阵乘优化的经典套路——分块tiling、把一块数据从HBM搬到共享内存、在寄存器里反复复用——本质上就是在这三级组织上做数据搬运的规划。我在实际写kernel的时候配置参数通常这么定block大小取128或256线程保证SM上有足够多的warp来掩盖访存延迟共享内存占用控制在SM可配置份额的一半以内给双缓冲留空间寄存器用量用-Xptxas -v看单个线程压到64个以内避免occupancy掉太多。这些数字不是教条但要有一个起点然后根据实测的occupancy和stall原因去调。2.3 NPU的指令粒度从MAC阵列到微码NPU的指令系统完全换了一套逻辑。它没有标量指令掩码这种表达指令的基本粒度是张量级的加载一个tile、做一个矩阵乘、写回一个结果块、激活函数、量化缩放。你用不着通常也不允许控制单个MAC单元在做什么那是硬件数据流自己决定的。为了让阵列一直忙NPU的数据通路一般是固定数据流形态权重固定在阵列里不动激活值从左侧按行推进部分和从上往下脉动传递。这叫脉动阵列systolic array好处是每个MAC单元只需要和相邻单元连线布线短、延迟低、能跑高频。代价是阵列尺寸固定比如128×128那么一次能处理的矩阵维度就被这个尺寸约束太小的矩阵喂不满阵列太大的矩阵要切分切分就带来额外的数据搬运。编译器在这里的角色比GPU上重得多。GPU上你写的代码和最终执行的指令之间隔着nvcc和PTXAS但语义基本保持一致NPU上你给的是一张计算图或者一段DSL描述编译器要做算子融合、内存分配、tile切分、流水排布、量化插入最后生成一长串微码。编译器做得好不好直接决定实测性能是标称的一半还是八成。我见过同一个模型在同一块NPU上只因为换了编译器的图优化开关端到端延迟差了将近40%。一个容易忽略的点NPU的不支持某算子很多时候不是硬件不支持而是编译器没有实现对应的拆解规则。遇到这种情况先翻编译器的支持列表和算子拆解文档再判断是不是真的硬件限制。2.4 从编译器视角看两种指令生成的差异把两边放到编译器视角对比一下差异更清楚。GPU编译器的任务相对轻它要做寄存器分配、指令调度、循环展开、地址强度削减但线程级并行是程序员显式给的编译器不需要去发明并行性。NPU编译器则必须在图上做并行性推导哪些算子可以融合成一个tile流水、中间结果放片上还是放DDR、双缓冲的深度设多少、量化参数在哪一层插入。这些决策的数量级远大于GPU编译器。这带来一个实际后果——NPU的性能调优更依赖工具链而不是手写代码的技巧。你在GPU上可以靠手写kernel榨出最后20%的性能在NPU上你基本只能在编译器的配置项里调参数剩下的交给它。所以选NPU方案的时候工具链成熟度这个软指标重要性不比峰值算力低。3. 数据通路拆解从HBM到片上SRAM的搬运账3.1 GPU的四级存储与带宽账本GPU的存储层级从上到下是寄存器堆、共享内存/L1、L2、HBM。以常见的数据中心加速卡为例寄存器堆容量在256KB/SM这个量级共享内存加L1合计192KB左右可配置L2在几十MB量级HBM带宽在1.5TB/s到3TB/s这个区间。注意这个层级里带宽的落差寄存器到ALU的带宽是每周期几百字节的量级共享内存每周期128字节左右L2大概几TB/sHBM落到1.5到3TB/s。也就是说从HBM到寄存器这条路带宽差了两个数量级以上。这就是为什么GPU优化里所有的手法归根到底都是一句话减少对HBM的访问次数。分块、复用、融合、预取全是为了这个。算一笔具体的账。假设一个算子要读入一个 4096×4096 的FP16矩阵数据量是32MB。如果HBM带宽按2TB/s算光读一遍就要16微秒。如果这个矩阵在一次计算里被复用了8次而没有做片上缓存那就变成128微秒而理论计算时间可能只有几微秒。带宽瓶颈就是这么出来的。3.2 NPU的暂存器与数据复用策略NPU走的是另一条路片上SRAM给得比较大几百KB到几MB量级而且带宽极高因为它是直接贴着计算阵列放的。很多NPU的架构图里SRAM和MAC阵列之间的总线宽度能做到每周期上千字节。这意味着它的带宽墙位置和GPU完全不同——HBM对它是慢速外存真正的性能取决于能不能把数据留在片上。所以NPU编译器的核心工作就是数据复用规划。以一次卷积为例输入特征图、权重、输出特征图三者里权重的复用价值最高每个权重会被用很多次输入特征图的复用取决于卷积核大小和步长输出特征图基本是一次性写入。数据流架构就是按这个规律分的权重固定数据流把权重常驻阵列激活值流动行固定数据流则反过来。选哪种取决于哪一类数据的复用次数更高。这里有个实际经验判断一个模型适不适合NPU不用看FLOPs看它的片上可容纳性。如果某层需要的中间张量远大于片上SRAM总量那么这层就必然要走DDR性能会明显掉下来。我一般会拿编译器的内存报告看每一层的片上占用峰值凡是超出SRAM容量的层都是重点优化对象通常通过切分batch或者改tile尺寸来解决。3.3 用Roofline算一遍算术强度想要一个不靠感觉的判断方法Roofline模型是最实用的工具。它的核心公式很简单可达到性能 min(峰值算力, 算术强度 × 带宽)。算术强度就是每读一字节数据能做多少次浮点运算。我写个简单脚本算一下def roofline(flops, bytes_moved, peak_tflops, bw_tbps): ai flops / bytes_moved # 算术强度 FLOP/Byte ridge peak_tflops * 1e12 / (bw_tbps * 1e12) # 平衡点 perf min(peak_tflops, ai * bw_tbps) return ai, ridge, perf # 大矩阵乘 MNK4096FP16 flops 2 * 4096**3 # 读 A、B写 C各 4096*4096*2 字节 bytes_moved 3 * 4096 * 4096 * 2 print(roofline(flops, bytes_moved, 300, 2.0)) # 输出算术强度约 1365 FLOP/Byte平衡点 150性能受算力限制 # elementwise add每元素读2写1算1次加法 flops2 4096 * 4096 bytes2 3 * 4096 * 4096 * 2 print(roofline(flops2, bytes2, 300, 2.0)) # 算术强度约 0.17性能被带宽死死限住只有 0.33 TFLOPS结果差距大得离谱同样是4096×4096的数据量矩阵乘能跑到算力峰值附近逐元素加法只能跑到峰值的千分之一。这解释了一个很常见的困惑——为什么有的算子优化半天毫无起色。因为它天生就是带宽受限的你就算把计算指令优化到极致性能也上不去唯一的出路是减少访存比如把前后两个逐元素算子融合成一个kernel中间结果不落HBM。同样的分析可以套到NPU上。假设某NPU的INT8算力是256 TOPS片上SRAM到阵列的带宽是5TB/s那么平衡点在51 OPS/Byte左右。因为片上带宽高、算力相对可控这个平衡点比GPU低不少意味着NPU在中等算术强度的算子上更容易跑满。这也是它在推理场景能效更好的一个底层原因跟NPU架构更先进没关系纯粹是数字关系的必然。4. 稀疏化结构化与非结构化是两条完全不同的路4.1 稀疏到底省了什么稀疏化的直觉很好理解既然权重大部分接近零那把这些零跳过去不算是不是就省了理论上省的是乘加次数但实际能不能省下来取决于硬件有没有能力跳过。这中间隔着一个巨大的工程鸿沟——跳过零意味着访存地址变得不规则而不规则访存恰好是硬件最讨厌的东西。先分清两类稀疏。非结构化稀疏是任意位置都可以为零压缩率高理论上限漂亮结构化稀疏要求零的分布符合某种规律比如每4个连续元素里恰好留2个非零这就是2:4稀疏。前者的压缩率能到90%以上但硬件很难加速后者的压缩率固定是50%但可以被硬件精确映射成更短的乘法流水线直接换来吞吐翻倍。我自己在做剪枝实验的时候有个体会非结构化稀疏带来的收益在通用GPU上几乎全部被索引开销和访存不连续性吃掉了。你确实少算了乘法但为了找到那些非零元素你多花了几倍的地址计算和加载指令净收益经常是负的。这就是为什么工业界最终选择了结构化方案——不是它理论更优而是它工程上可落地。4.2 2:4结构化稀疏在GPU上的落地细节2:4稀疏的硬件支持方式值得仔细看一下。做法是在每4个连续元素里只保留绝对值最大的2个另外2个置零同时记录被保留元素的索引通常2个bit。硬件在做矩阵乘的时候读入压缩后的数据和索引用一个选择网络把另一侧的激活值挑出对应位置然后送进缩短了的乘法阵列。因为乘法器数量只有一半在用理论上吞吐就是2倍。一个典型的剪枝流程大致长这样import torch def apply_2_4_mask(weight): # weight: [out_features, in_features] w weight.clone() # 按每4个一组切分 grouped w.view(w.shape[0], -1, 4) # 找出每组里绝对值最小的2个 _, idx grouped.abs().topk(2, dim-1, largestFalse) mask torch.ones_like(grouped, dtypetorch.bool) mask.scatter_(-1, idx, False) return (grouped * mask).view_as(w) # 注意这不是一次剪枝就完事之后必须做微调恢复精度流程是四步先正常训练到收敛然后施加2:4掩码接着用较小的学习率微调若干轮让精度回补最后导出时确认框架真的启用了稀疏kernel。第三步不能省我试过直接剪完就部署精度掉了两个多百分点加了一轮短微调之后基本追平了稠密模型。这里有个容易踩的坑光把权重置零不等于启用了稀疏加速。你得确认推理框架真的把权重转成了压缩格式、真的调用了稀疏版本的矩阵乘kernel。有些框架默认只是把零保留在稠密张量里那你的收益是零还得额外付出剪枝带来的精度损失。验证方法很直接对比稠密和稀疏两条路径的实测延迟如果差不到20%基本可以判断稀疏路径没被真正启用。4.3 NPU对稀疏的支持差异与验收陷阱NPU这边情况更复杂一些因为不同厂商支持的稀疏形态差别很大。有的只支持结构化的块稀疏比如4×4的小块整体剪掉有的支持权重稀疏但不支持激活稀疏有的在编译器层面提供了稀疏通路但实际kernel还没接通。这些差异在参数表上通常看不出来必须在实测里确认。验收的时候我会做三件事。第一拿一个确定能被稀疏化的层比如某层的大矩阵乘单独测稠密和稀疏两条路径看延迟比是不是接近理论值。第二把稀疏模型和稠密模型跑同一个测试集比较精度差异是否在可接受范围内。第三看编译器生成的报告里稀疏相关的标记是否真的存在。三步都过才算稀疏这条路真的通了。说个反直觉的现象稀疏化在训练场景里通常收益比推理小得多因为反向传播的梯度往往是稠密的剪枝带来的加速会被梯度计算的开销抵消。如果你的目标只是推理加速那就专注于推理侧的稀疏通路别指望训练也能跟着一起快。5. 存算一体把乘法搬进存储阵列里面5.1 冯诺依曼瓶颈到底卡在哪前面反复说的带宽墙本质是冯诺依曼结构的老问题计算单元和存储单元物理分离数据必须在两者之间来回搬。搬到今天搬数据的能耗已经远超计算本身的能耗——在很多工艺节点的实测里一次HBM访问的能耗是一次FP32乘加的几十倍甚至上百倍。这意味着即使你把算力堆到天上只要数据还得从HBM搬过来能效比就被锁死了。存算一体Compute-in-MemoryCIM的思路就是把计算挪进存储阵列里让数据不动、计算动。最典型的实现是阻变存储器RRAM交叉阵列每条行线和列线的交叉点放一个阻变器件电导值代表权重输入电压加在行线上根据欧姆定律流过每个器件的电流等于电压乘电导同一列上所有电流按基尔霍夫电流定律自然累加得到的列电流就是这一列的乘加结果。整个矩阵乘在一次读操作里就完成了不需要把权重搬出来。这个方案的诱惑力在于它把搬运这件事彻底消掉了理论上能效比可以提高一到两个数量级。但真实的工程难度也在这里——模拟域的乘法对器件特性极其敏感。5.2 CIM阵列的电路原理与精度补偿用RRAM做CIM要面对至少四类非理想因素。器件电导的写噪声和漂移会让权重值随时间变化今天写进去的矩阵明天读出来可能就偏了阵列里的线电阻造成IR drop离驱动端越远的单元实际电压越低ADC和DAC的精度和功耗往往成为整个方案的瓶颈因为把模拟电流转成数字需要高精度比较器而高精度又意味着高功耗和大面积还有器件良率和一致性阵列越大良率问题越突出。实践中能做的补偿手段有几类。写验证write-verify通过多次写入和回读校准把权重误差压到可接受范围内差分对结构用两个器件表示一个权重正负分开抵消共模漂移把模拟计算限制在低位宽比如4到6 bit超出部分用数字路径补齐在算法侧做量化感知训练让模型对权重误差更鲁棒。这些手段加起来目前的CIM方案大致能做到4到8 bit的有效精度适合推理里对精度不那么敏感的层比如某些卷积层对精度敏感的层比如归一化、注意力里的softmax附近还是得放回数字路径。所以现在几乎没有纯CIM的方案都是模拟域做一部分、数字域做一部分的混合架构。5.3 现阶段可用的几种折中形态如果你的目标是现在就能用上的技术那纯CIM基本还停留在研究和早期产品阶段真正可用的是几类折中方案。第一类是近存计算PIM把计算单元放在存储芯片内部或者紧贴存储控制器缩短搬运距离但不改变数字计算的本质比如在HBM的每个die里塞入若干计算单元这样可以把带宽压力从片外转到片内。第二类是权重驻留方案把模型权重常驻在加速卡的片上或者近存区域激活值流动这其实就是很多NPU在做的事只是表述不同。第三类是把CIM用在最规则的运算上——比如推荐系统里的embedding查表——这类运算的访问模式极其规则对精度要求也相对宽松很适合做验证。我个人的判断是未来几年最现实的路径不是完全替换而是分层极端规则的运算下沉到存内中等复杂度的走NPU的数据流阵列剩下的通用部分留给GPU。三层各干各的活谁也别想通吃。6. 常见问题与排查技巧实录6.1 一张速查表覆盖八成现场问题实际操作里遇到的问题高度重复我整理了一张表按现象直接查原因能省掉很多试错时间。现象可能原因排查动作GPU上算力利用率长期低于30%算子带宽受限或kernel launch开销过大用profiler看memory throughput和compute throughput哪个先打满显存占用不高但程序卡住同步阻塞、CPU侧等待或数据加载瓶颈检查是否在循环里做了逐样本的同步拷贝NPU上某个算子回退到CPU编译器不支持该算子拆解查看编译日志里的fallback记录NPU精度比GPU低明显量化参数标定不足或某个层精度敏感逐层对比输出定位偏差最大的层稀疏模型没有加速稀疏kernel未启用权重仍是稠密存储对比稠密/稀疏两条路径的实测延迟多卡扩展效率低通信和计算没有重叠检查是否用了异步通信、梯度分桶是否合理NPU端到端延迟远高于纯推理时间前处理和后处理跑在CPU上把图像resize、NMS等搬到加速器模型转换后性能骤降图优化阶段丢失了算子融合机会对比转换前后计算图结构变化6.2 我踩过的几个坑第一个坑是关于显存看起来够用的误判。我遇到过显存占用显示只有一半但训练一直在卡的情况。后来发现是碎片化反复申请释放不同尺寸的张量之后空闲显存被切得七零八落大块分配失败。解决办法是提前规划好缓冲区或者开启内存池的碎片整理策略。这个问题在长时间运行的服务里特别常见。第二个坑是量化标定的数据分布。做INT8量化的时候我用了一个小批次的校准集结果模型在某类样本上精度崩了。原因是校准集的分布和真实流量不匹配那些没见过的数值范围被截断了。后来我把校准集扩到覆盖各种极端场景并且对激活值用了逐通道的缩放系数问题才解决。经验就是量化标定的数据宁可多花时间采样也别图省事用现成的。第三个坑是把峰值算力当成实际能力。选型阶段我按参数表上的TOPS做容量规划结果实测只跑到三成。回头复盘那个标称值是在理想batch、理想shape、全部数据在片上复用的条件下测出来的我的实际业务里batch很小、序列长度变化大阵列根本喂不满。后来我改成用实测的每秒处理样本数来做容量规划再留30%余量上线就再也没出过容量事故。第四个坑是关于指令集兼容性的。有次我按文档写了一个自定义算子在某个固件版本上跑得好好的升级固件之后直接编译失败。原因是新版本的编译器调整了内部指令的编码和约束。教训是自定义算子一定要锁版本别跟着固件随便升。如果非要升先在测试环境跑完整的算子回归。6.3 选型时我会问自己的四个问题最后分享我实际做判断时会走的四步。先问模型结构稳不稳定——如果算子集合每周都在变通用性优先选GPU如果结构冻结了就往下走。再问瓶颈在哪——用Roofline粗算一遍主要的几层如果大部分层算术强度很高NPU有机会如果一堆逐元素操作主导带宽是主要矛盾那要看谁的片上复用做得好。然后问工具链成熟度——编译器的算子覆盖率、调试工具的完善程度、社区资料的多少这些软指标会直接决定你的迁移成本。最后问精度底线——能不能接受INT8甚至更低有没有敏感的层必须保留高精度稀疏化之后的精度回补周期能不能接受。这四个问题问完答案通常就出来了不太需要再纠结参数表上那几个数字。我个人在实际项目里的体会是GPU和NPU的关系更像钳子和扳手而不是新老交替。真正把项目做顺的人往往是对两边都摸过底的人——知道GPU的warp为什么发散、知道NPU的阵列为什么喂不满、知道稀疏的收益卡在哪一层、知道存算一体现在的边界在哪。这些知识不会写进参数表但它们决定了你手里的硬件到底能发挥出几成。