GPU与NPU推理加速选型：指令集、数据通路、稀疏化与存算一体

手里同时压着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的核心执行模型是SIMT（Single 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、block（CTA）、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/s，HBM落到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 # 大矩阵乘 M=N=K=4096，FP16 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, largest=False) mask = torch.ones_like(grouped, dtype=torch.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-Memory，CIM）的思路就是把计算挪进存储阵列里，让数据不动、计算动。最典型的实现是阻变存储器（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 一张速查表覆盖八成现场问题

实际操作里遇到的问题高度重复，我整理了一张表，按现象直接查原因，能省掉很多试错时间。

6.2 我踩过的几个坑

第一个坑是关于显存"看起来够用"的误判。我遇到过显存占用显示只有一半，但训练一直在卡的情况。后来发现是碎片化：反复申请释放不同尺寸的张量之后，空闲显存被切得七零八落，大块分配失败。解决办法是提前规划好缓冲区，或者开启内存池的碎片整理策略。这个问题在长时间运行的服务里特别常见。

第二个坑是量化标定的数据分布。做INT8量化的时候，我用了一个小批次的校准集，结果模型在某类样本上精度崩了。原因是校准集的分布和真实流量不匹配，那些"没见过"的数值范围被截断了。后来我把校准集扩到覆盖各种极端场景，并且对激活值用了逐通道的缩放系数，问题才解决。经验就是：量化标定的数据，宁可多花时间采样，也别图省事用现成的。

第三个坑是把峰值算力当成实际能力。选型阶段我按参数表上的TOPS做容量规划，结果实测只跑到三成。回头复盘，那个标称值是在理想batch、理想shape、全部数据在片上复用的条件下测出来的，我的实际业务里batch很小、序列长度变化大，阵列根本喂不满。后来我改成用实测的每秒处理样本数来做容量规划，再留30%余量，上线就再也没出过容量事故。

第四个坑是关于指令集兼容性的。有次我按文档写了一个自定义算子，在某个固件版本上跑得好好的，升级固件之后直接编译失败。原因是新版本的编译器调整了内部指令的编码和约束。教训是：自定义算子一定要锁版本，别跟着固件随便升。如果非要升，先在测试环境跑完整的算子回归。

6.3 选型时我会问自己的四个问题

最后分享我实际做判断时会走的四步。先问模型结构稳不稳定——如果算子集合每周都在变，通用性优先，选GPU；如果结构冻结了，就往下走。再问瓶颈在哪——用Roofline粗算一遍主要的几层，如果大部分层算术强度很高，NPU有机会；如果一堆逐元素操作主导，带宽是主要矛盾，那要看谁的片上复用做得好。然后问工具链成熟度——编译器的算子覆盖率、调试工具的完善程度、社区资料的多少，这些软指标会直接决定你的迁移成本。最后问精度底线——能不能接受INT8甚至更低，有没有敏感的层必须保留高精度，稀疏化之后的精度回补周期能不能接受。

这四个问题问完，答案通常就出来了，不太需要再纠结参数表上那几个数字。

我个人在实际项目里的体会是，GPU和NPU的关系更像钳子和扳手，而不是新老交替。真正把项目做顺的人，往往是对两边都摸过底的人——知道GPU的warp为什么发散、知道NPU的阵列为什么喂不满、知道稀疏的收益卡在哪一层、知道存算一体现在的边界在哪。这些知识不会写进参数表，但它们决定了你手里的硬件到底能发挥出几成。