★ FEATURED ARTICLE 第一次拿到 SambaNova 芯片的架构资料时我盯着那张框图看了挺久——上面没有 GPU 那种熟悉的 L2 缓存和 SIMT 调度器取而代之的是一大片重复排列的 PCU 和 PMU中间靠一张多级片上网络串起来。这种布局在通用加速器里不常见它更像把 FPGA 的灵活性削掉一部分换成更粗的粒度再把省下来的面积全砸到数据通路上。这颗芯片走的路线叫可重构数据流架构核心单元叫RDU也就是 Reconfigurable Dataflow Unit。它想解决的问题很直接在 AI 负载里算力早就不是瓶颈了把数据从内存搬到计算单元所花的能量和时间才是。围绕这一点我把这套架构从头到尾捋了一遍包括硬件怎么划分、编译器在背后做了什么、算力和带宽怎么配平、真上手要踩哪些坑下面全部分享出来。适合正在做推理集群选型的人、想搞清楚数据流架构和 GPU 到底差在哪的人也适合刚接触异构加速、想补一节架构课的朋友。1. 先弄清楚这颗芯片想解决什么问题1.1 数据搬运才是真正的成本要理解任何一颗 AI 加速器第一步不是看它有多少个计算单元而是看它的钱花在哪。业界常引用的那组经典数据是一次 32 位浮点加法大约消耗 0.9 pJ从片上 SRAM 读一次约 5 pJ从片外 DRAM 读一次约 640 pJ。也就是说一次片外访存的能量开销够你做七百次加法。这个数量级差距直接决定架构设计的方向。GPU 的做法是尽量把数据往寄存器堆和共享内存里塞靠大量线程的切换把访存延迟藏起来代价是要维护庞大的缓存层次以及一个相当复杂的调度器。数据流架构的思路则是另一个方向——干脆不让数据来回跑让计算单元按固定拓扑连成流水线数据像流水线一样从上游流到下游流过一遍就算完。这种计算单元固定、数据流动的思路正是可重构数据流的核心也是 SambaNova 芯片敢把缓存层次砍掉的底气所在。问题在于AI 模型的拓扑是变化的。今天的卷积网明天就换成 Transformer如果做一颗纯固定流水线的 ASIC等于赌死一种结构风险太大。所以折中方案才是主流做一批可以被重新布线的粗粒度计算单元每次跑新模型之前用编译器把模型的图映射上去物理上连成一条新的流水线。这就是可重构三个字的真实含义——它不是运行时动态切换而是在编译期就定义好数据通路。1.2 三条路线的分野时间复用、空间展开、固定流水线把当下主流的几类架构摆在一起对比能看得很清楚。架构类型执行方式灵活性单位算力能效编程模型GPUSIMT时间复用同一批计算单元轮流跑不同算子高中成熟生态kernel 编程FPGA位级可重构空间展开极高低到中HDL 或高层次综合门槛高固定流水线 ASIC空间展开硬件写死无最高基本不可编程可重构数据流本芯片空间展开编译期重配置中高张量框架图直接下沉区别的关键在于空间和时间这两个维度。GPU 属于典型的时间复用它有几千个核心但同一个核心在不同时刻执行的是不同算子算子之间来回换每次切换都要重新加载权重、重新写回中间结果。而可重构数据流是把空间维度用满——同一时刻芯片上不同区域跑的是同一个模型的不同层前一层的输出直接流进后一层中间结果根本不需要落到片外内存。这就是为什么这类架构在小 batch 推理上能打出优势。中间结果不落片外内存带来的不只是省带宽还省掉了写回、同步、重新读取这一整套开销。而代价同样明显一条流水线一旦搭起来它就被这次输入的尺寸固定住了输入 shape 一变化流水线可能得重新编。这个矛盾贯穿了后面所有的工程问题。2. 硬件架构拆解三种计算单元与一张片上网络2.1 PCU粗粒度可重构的计算格子PCU 全称 Pattern Compute Unit是整个片子的算力来源。从公开的架构描述看单个 PCU 内部并不是一堆同质的 ALU而是一个小车间里面有 SIMD 的向量 ALU 处理常规的乘加有独立的浮点单元处理精度要求更高的部分还有若干可定制的功能单元能做非线性激活、比较、位运算这类固定模式的操作。这些单元之间通过一个交叉开关crossbar连到统一的寄存器堆上。这种设计的用意很实际。如果每个 PCU 都做成通用处理器那和 GPU 没区别如果做成纯乘加阵列那激活函数、归一化这些非矩阵运算就得丢回主机或额外的单元去算流水线立刻断掉。粗粒度可重构的意思是粒子上我保留足够的通用性但这个通用性是在编译期确定的不需要为运行时判断分支预留开销。实测经验里PCU 的结构对这类模型特别友好——它的运算密度高几乎每个操作都是规整的向量运算可以让 PCU 长时间满负荷不需要靠多线程切换去填延迟。相比之下控制流多、分支多的模型在 PCU 上就很难受原因是它没有乱序执行、没有深度多线程一旦执行流停下来整个 PCU 就在空转。注意判断一个模型适不适合这类芯片第一件事是看它的控制流密度而不是看它的 FLOPs。分支密集的模型即便算力需求再大也不划算。2.2 PMU 与 AGU把访存变成本地事务PMU 是 Pattern Memory Unit本质上是带地址生成逻辑的片上暂存器scratchpad。它和缓存最大的区别在于缓存是靠硬件猜你下一步要什么命中率不确定暂存器是编译期就规划好哪块数据放哪完全确定。AGU 则是 Address Generation Unit专门管地址计算和数据搬运。传统处理器里地址计算要占用宝贵的指令槽和寄存器把它们剥离出来交给独立的 AGU让 PCU 只关心算术整体效率会好不少。这套设计让片外访存变成了一个可以被调度的动作——编译器明确知道什么时候要搬什么可以提前发起搬运让计算和加载重叠。这里有个容易被忽略的细节片上 SRAM 的容量。这颗芯片公开资料里给出的量级是百 MB 级别具体数字随代际变化以官方文档为准。百 MB 听起来不少但要装下一个十几 GB 的大模型权重显然不可能。所以真正的工程做法是分时复用——把模型按层切块每一时刻只有当前活跃的几层在片上其余权重留在片外靠高带宽接口按流水节奏喂进来。2.3 片上网络与多芯片扩展单个芯片内部的 PCU 和 PMU 之间靠一张多级层次的片上网络连接。这张网络的设计目标不是低延迟而是高吞吐和可预测性。因为在数据流执行里每条边上的数据速率在编译期就确定了网络只需要按这个速率稳定地送数据不需要应对随机的、突发式的访问。往外扩展时多颗芯片之间走的是基于标准以太网物理层的定制互连。这个选择挺有意思——不自己造一套私有物理层直接用成熟的以太网 SerDes成本和风险都低。在软件层面多芯片被抽象成一个统一的计算资源池编译器负责把一条大流水线切分到多颗芯片上跨芯片的边就变成网络传输。这时候模型并行、流水线并行的切分策略就成了决定扩展效率的关键。组件职责设计取向工程上的影响PCU向量/矩阵运算、激活粗粒度可重构编译期定功能控制流密集的算子效率低PMU片上暂存、查找表确定性寻址不做缓存替换容量规划必须精确超了要重编AGU地址计算与 DMA计算与搬运解耦访存可预取掩盖延迟片上网络单元间数据传递高吞吐、速率可预测拓扑变化触发重新编译芯片间互连多卡扩展基于标准以太网物理层大模型切分策略直接影响效率3. 高性能的真正来源把编译器当硬件用3.1 空间编译一次配置多级流水同时跑如果要用一句话概括这套架构的高性能秘诀那就是硬件的行为在编译期就全定了运行时几乎没有决策开销。GPU 的执行模型是取指—译码—调度—执行每一轮循环都有指令开销即便有指令缓存和发射优化指令带宽依然占用了可观的资源和功耗。数据流架构把这一整层都省掉了编译器把计算图切成若干阶段每个阶段绑到一组 PCU 上阶段之间用片上网络连成流水配置完成后硬件就一条道跑到黑。举个具体例子。一个三层的块结构# 伪代码一个可融合的块用来说明图切分 def block(x, w1, w2): h matmul(x, w1) # 阶段 0 h gelu(h) # 阶段 1 y matmul(h, w2) # 阶段 2 return y在 GPU 上这三步大概率是三个 kernel中间结果 h 要写回显存再读出来两次访存就是两次开销。在数据流架构上编译器会把这三步分别映射到三组 PCU中间插两级 FIFO 缓冲数据从阶段 0 出来直接进阶段 1全程不落片外。同时因为三个阶段在不同硬件上并行吞吐是可以叠加的——理想的完全流水情况下整体吞吐等于最慢那一级的速度而不是三级之和。这就是空间展开带来的本质差别GPU 是三步串行做完数据流是三步同时做不同批的数据。当阶段数足够多这个差距会非常可观。3.2 算子融合与缓冲插入的实际做法融合不是免费的。把两个阶段连起来前提是两边的数据速率能对上——如果阶段 0 每周期产出 128 个元素阶段 1 每周期只能消耗 64 个那中间就必须插缓冲而缓冲要占 PMU 的容量。所以编译器的核心工作其实是三道题第一道是切分。把计算图切成粒度合适的阶段太粗了硬件放不下太细了流水线开销占比高。这个粒度通常以单个 PCU 或一组 PCU 的承载能力为单位。第二道是负载均衡。理想情况下每级耗时接近流水线才不会出现空转。编译器会做重新的算子拆分比如把一个大矩阵乘拆成两半放到两组 PCU 上让它们并排跑。第三道是缓冲分配。两个阶段速率不匹配时要么插入 FIFO 做背压要么把上游数据在 PMU 里积攒成大块再一次性送下去。这一步直接决定片上 SRAM 的占用率占满了就得往片外溢性能立刻断崖。# 伪代码片上容量预算的粗算用来判断切块是否可行 def check_tile(h, w, keep, elem_bytes2, total_sram_mb200, safety0.7): # keep: [batch, seq, hidden] 这类单层的中间激活 act keep[0] * keep[1] * keep[2] * elem_bytes # 权重块 激活 双缓冲粗略按 3 倍算 need act * 3 budget total_sram_mb * 1024 * 1024 * safety return need budget, need / (1024 * 1024)这类估算在真上手之前一定要跑一遍。我见过不少团队一上来就把大模型整层往上怼结果编译直接报容量超限浪费两三天才回到正确路径。3.3 编译时间为什么长以及怎么压这是这类平台最常被吐槽的一点编译动辄几分钟到几十分钟比 GPU 上编译一个 kernel 慢了一个数量级。原因不神秘。GPU 的编译器只需要把单个算子编译成一段 SASS算子之间是独立的而数据流编译器要做的是一次全局布局——把整张图映射到一个有限的、带拓扑约束的资源池上这是一个带约束的优化问题。图越大搜索空间越大。压缩编译时间的经验有几条。一是固定 shape别让每批数据都触发重新编译把可变长度 pad 到固定桶里。二是分层编译并缓存只重编改动的子图没改的部分复用产物。三是提前预编译热门配置把常见 batch size 的组合在离线阶段全都编好线上直接加载。注意编译产物一定要版本化管理。模型权重、编译器版本、编译配置三样任意一个变了产物都得重出。我踩过一次坑线上环境用了旧产物输出结果对不上排查了大半天才发现是版本漂移。4. 用 roofline 算一遍算力和带宽的对账单4.1 算术强度门槛的推导任何加速器性能上限都能用 roofline 模型框出来。核心是两个量算力 PFLOP/s和带宽 BByte/s它们的比值 P/B 就是算术强度门槛。下面用一组示意参数来演示算法不代表任何实测数字当真值请以官方文档为准。假设某代芯片的 BF16 算力 P 在数百 TFLOPS 量级取 300 TFLOPS片外带宽 B 取 1 TB/s。那么门槛就是300e12 FLOP/s ÷ 1e12 Byte/s 300 FLOP/Byte意思是一个算子每从片外搬 1 字节数据至少要完成 300 次浮点运算才能跑到算力上限做不到就是带宽受限。4.2 大 GEMM 与小 batch decode 的两种命运拿一个大矩阵乘来算。当 M N K 4096 时计算量 2 × M × N × K 1.37e11 FLOP约 137 GFLOP数据量 (M×K K×N M×N) × 2 字节 100 MB算术强度 1.37e11 ÷ 1.0e8 ≈ 1370 FLOP/Byte远超门槛妥妥的算力受限性能不会浪费。这也是为什么离线训练和 prefill 阶段在这类芯片上表现好。再看大模型推理的 decode 阶段。此时每个 token 的生成本质上是一个向量乘矩阵M 1N、K 是隐藏维度比如都是 4096。计算量 2 × 1 × 4096 × 4096 ≈ 3.4e7 FLOP数据量 (1×4096 4096×4096 1×4096) × 2 ≈ 33.6 MB算术强度 ≈ 1 FLOP/Byte比门槛低了两个数量级以上。这个阶段是纯纯的带宽受限算力和理论峰值一点关系都没有实际性能完全由每秒能搬多少权重决定。4.3 权重驻留decode 场景的破局点带宽受限这件事没法靠堆算力解决能改的只有两条路。第一条是增大 batch。把 M 从 1 提到 32计算量涨 32 倍权重数据量不变算术强度直接提到 32 FLOP/Byte。虽然还是不到门槛但已经好了一个半数量级。这是最直接的优化手段代价是显存占用和首字延迟。第二条是权重驻留。既然 decode 的问题是把权重一层层从片外读进来那能不能反过来——把权重按层切块后钉在片上 SRAM 里让输入数据流过去这正是数据流架构擅长的模式权重一次性加载进 PMU同一块权重服务整批请求数据流过就出结果。片上 SRAM 有百 MB 量级装不下整模型但装几层没问题配合层间流水可以做到上一层还在算下一层的权重已经在加载。# 伪代码层间流水的节奏控制说明权重预取的思想 def pipelined_decode(layers, x, budget_mb): prefetch(layers[0].weight) # 先把第 0 层搬进来 for i, layer in enumerate(layers): if i 1 len(layers): prefetch(layers[i 1].weight) # 算第 i 层时预取第 i1 层 x layer.forward(x) # 计算与搬运重叠 return x不过这里有个现实约束如果模型并行切分做得不好片上 SRAM 时刻装着一堆用不上的数据效果就打折。所以片上容量规划不是能塞多少塞多少而是要和流水线节奏严格对齐。5. 落地实操从模型到可执行流水的完整链路5.1 环境接入与模型迁移的第一步拿到一台这样的机器接入流程和 GPU 集群差别不小。典型路径是驱动与运行时安装 → 编译器接入 → 用张量框架写模型 → 编译器抓图 → 编译产物 → 执行。整个链路里最关键的是抓图这一步因为编译器要拿到完整的计算图才能做全局优化。因此第一个实操动作是把模型改写成可被完整捕获的形式。这意味着动态 shape、动态控制流要提前消解掉尽量转成固定形状Python 侧的数据预处理逻辑不能混在图里要拆到前处理阶段自定义算子要用框架提供的可捕获方式实现普通 Python 循环走不通。大部分团队卡在第一步原因往往不是芯片不好用而是模型本身写得太Pythonic有大量的动态分支和算子外逻辑。我的建议是先在 CPU 上把模型改造成静态图能跑通、结果对齐再上加速卡能省掉一半的排查成本。5.2 一个 GEMM 加激活的映射实例一个最简单的映射就拿矩阵乘加激活来看。整个过程的思考顺序是这样的先决定分块。用一个静态的切分策略把大矩阵切成能被单组 PCU 一块吃完的尺寸def pick_tile(M, N, K, sram_bytes, elem_bytes2, safety0.7): budget sram_bytes * safety # 留给单个阶段的容量 best None for tm in (32, 64, 128, 256): for tn in (32, 64, 128, 256): for tk in (16, 32, 64, 128): a tm * tk * elem_bytes # A 块 b tk * tn * elem_bytes # B 块 c tm * tn * 4 # C 块用 FP32 累加 if a b c budget and 2 * tm * tn * tk 2 * 64 * 64 * 64: best (tm, tn, tk) return best这段代码的思路是在容量约束下尽量让单块的计算量够大因为块越小计算和访存的比例越差。2*tm*tn*tk是这一块的浮点运算量右边那个判定是给一个下限防止选出一个算力根本喂不饱的块。再分配阶段。把GEMM放到一组 PCU把激活放到另一组 PCU两级之间插 FIFO。此时要注意激活的输入输出形状要匹配tm × tn的块输出直接送过去别做转置、别做 reshape否则编译器会插一个额外的搬运阶段白白吃掉带宽。最后压性能。三个可调的旋钮分块大小、流水级数、片上缓冲深度。这三者互相拉扯通常的做法是先固定分块扫流水级数找到吞吐拐点再回头微调分块。5.3 调优的三个抓手与参数取值实际调优时我一般按这个顺序动。调优抓手调整方向观察指标常见取值区间分块大小增大到容量边缘再回退单阶段吞吐、SRAM 占用率128~512 元素级别流水级数逐步增加直到吞吐不再上升单个 batch 的端到端时延2~8 级批量大小在时延约束内尽量加大单位时间 token 数由 SLA 反推缓冲深度按上下游速率差设置流水线气泡率1~4 个缓冲槽精度组合BF16 计算 FP32 累加输出与基准的数值偏差视模型敏感度这里有个反直觉的点批量大小不是越大越好。流水线级数固定时batch 太大反而会让每一级排队变长端到端时延飙升而有些线上场景对首字延迟有硬要求那就必须牺牲吞吐换时延。这个权衡在 GPU 上通过 batch 调度就能做在这类平台上还得额外考虑编译期确定的分块尺寸是否兼容。6. 踩坑实录问题现象、排查路径与速查表6.1 五类高频问题的现场记录编译超时或卡死。最常见的触发原因是计算图太大或者图里藏着控制流。排查时先把图导出来看规模超过一定节点数就分块编译把大模型拆成几个子图分别处理控制流要么改成静态掩码要么挪到图外。第一轮推理慢得离谱。这通常不是性能问题而是编译和加载开销被打进了计时。很多框架的第一次调用会触发即时编译包含编译、加载产物、初始化内存几件事几十秒很正常。解决办法是预热上线前跑几十次空转把产物缓存住。吞吐只有预期的一半不到。优先怀疑流水线没填满。检查批量大小是不是太小、流水级数是不是设得太浅、上下游速率是否不匹配导致气泡。用 profiling 看每一级的利用率找最低的那一级它就是瓶颈。性能随输入长度剧烈波动。十有八九是 shape 变化触发了重新编译。解决方案是把输入 pad 到固定桶或者预编译一组常用长度。这个坑非常隐蔽因为从业务侧看只是有时候快有时候慢。数值结果和 GPU 对不上。先分级排查小规模矩阵乘对不上是算子实现或累加精度的问题大规模对不上是并行切分里累加顺序差异导致的浮点误差累积只有边界元素对不上那大概率是 padding 的处理方式不一致。6.2 常见问题速查表现象最可能原因排查动作处理手段编译卡在某个阶段图过大或含动态控制流导出中间表示看节点数拆子图、消解分支首轮时延异常高即时编译与初始化开销打点看耗时分布预热、预编译产物吞吐低于预期一半流水线未填满看各级利用率加批量、调流水级数报容量超限单片 SRAM 放不下跑容量预算脚本缩小分块、减少驻留层数性能抖动明显输入 shape 触发重编统计 shape 分布pad 到固定桶精度下降累加精度不足或量化误差与 FP32 基准逐层比对改 FP32 累加、跳过敏感层多卡扩展效率低切分策略导致通信占比高测通信与计算时间占比调整切分点减少跨卡边6.3 几条没法从文档里学到的经验第一条先做小模型打通全链路别一上手就怼大模型。我见过太多团队直接上几十亿参数的模型编译报错、容量超限、精度不对全堆在一起根本分不清是哪一环的问题。正确的顺序是先用一个小到能一眼看穿结果的模型比如两三层、单头注意力跑通全链路确认数据流入流出、精度对齐再逐级放大。第二条把片上容量当成稀缺资源去规划而不是当成够用就行。在 GPU 上SRAM 不够了就溢出到显存性能降一点但还能跑在这类平台上容量规划失败往往直接编译不过或者流水线结构被迫变形性能不是降一点而是断崖。所以每次改模型第一件事就是重跑一遍容量预算。第三条不要指望把 GPU 上的调优直觉直接搬过来。GPU 上调优靠的是 occupancy、线程块大小、共享内存 bank 冲突这套东西在这类平台上等价的概念变成了分块尺寸、流水级数、缓冲深度。拿着 GPU 的经验去调经常会调反方向。提示多卡场景下优先把跨卡通信的边做到最少而不是去优化单卡的算力利用率。跨卡那一段的带宽通常比片上低一个数量级以上一条跨卡边就足以把整体流水线拖住。7. 选型边界什么活该交给它什么活别硬塞7.1 适配场景的判断清单不是所有负载都适合这类架构。我总结了一个粗筛清单能快速判断一个模型值不值得往这条路上走计算图是否相对静态如果每次推理都走不同的分支编译期没法确定流水线那就不合适。输入 shape 是否稳定如果长度分布极散且无法分桶重编译成本会吃掉所有收益。算子是否以矩阵乘、卷积、归一化、激活为主这几类能被高效映射稀疏、图结构、复杂索引类的算子很难。批量是否能做起来在线场景如果只能 batch1收益会打不少折扣。对首字时延的容忍度如何流水线架构的时延和吞吐是反比关系两者都要极致很难。按这个清单过一遍能筛掉大半不合适的场景省下很多无用的验证时间。7.2 评估一个数据流平台该看什么如果正在做选型除了看峰值算力这个数字参考价值其实有限我建议重点看四件事。一是编译器的成熟度。这是整套方案的核心编译器不行硬件再好也发挥不出来。具体看常见模型能不能开箱编译通过、编译时间是否可接受、报错信息是否可读。报错信息可读性这一条特别容易被忽略但它直接决定了排查效率。二是算子库的覆盖度。百分百覆盖不现实重点看你自己的模型里那十几个关键算子是否原生支持。不支持的算子要么能自动回退要么能用自定义方式实现两者都没有就是硬伤。三是调试工具链。出了问题能不能看到每一级的利用率、缓冲区水位、通信耗时看不到的话性能调优基本靠猜这在生产环境里是不可接受的。四是精度验证的便利性。能不能方便地和基准结果逐层比对如果只能比最终输出那定位精度问题的成本会高得离谱。我个人的经验是评估阶段花在工具链上的时间和后期上线省下的时间是成正比的。硬件参数差距百分之二三十可能一次调优就追回来了工具链差一代那就是数月的效率差距。最后说个实操体会这类可重构数据流架构真正的门槛不在硬件而在你得学会用编译器的视角去写模型。同样一个网络写法上稍微调整一下——把动态逻辑移到图外、把 shape 固定下来、把算子换成编译器认识的等价形式——性能可能差出一倍。刚接触的时候我总觉得这是硬件在逼着人迁就它后来发现这个约束其实逼着你把模型结构想得更清楚反过来对模型本身的优化也有帮助。
阅读完成 · 觉得有帮助？
