IREE编译器：如何将调度与执行一起烧进产物并压到10KB

最近几年我做模型部署相关的工作，接触到 IREE 编译器之后，很多之前绕不过去的难题确实换了个解法。其中一个很具体的痛点：有些嵌入式场景闪存是按 KB 算的，而传统方案跑一个模型要带"一个解释器 + 一份图描述 + 一堆 kernel 库"，这套东西体积压不下来，调度逻辑还只能留在大而全的运行时里，改一处调度策略就得重新烧整个固件。IREE 的思路是把调度和执行一起编译进一个产物——理想情况下可以压到 10 KB 级别——再交给一个几乎没啥逻辑的迷你运行时去执行。这篇文章不打算泛泛吹 IREE 多厉害，我按"为什么要这么设计 → 编译链路上调度和执行分别在哪儿成形 → 10 KB 怎么抠出来的 → 实操怎么复现 → 我踩过什么坑"这条线，把它从调度到执行的核心链路讲透。适合正在做模型部署选型的同学、想了解编译器如何参与 AI 推理的同学，以及准备评估 IREE 的架构师。

1. 调度要不要写进二进制里：我在这件事上吃的教训

先说个背景故事。之前我维护一个带循环结构的模型，输入端不定长，内部要跑 while 循环。用传统图解释器部署到 MCU 上时，循环的"调度逻辑"是运行时解释器的一部分——也就是解释器程序里去遍历子图、判断条件跳转、管理循环栈。模型一旦复杂点，解释器本身要增加功能，整个 runtime 固件都要跟着升级。最尴尬的是 kernel 库和图描述还是分开部署的，两边版本对不上时，报错信息基本等于没有，全靠肉眼排查。

那会儿我就意识到一个问题：调度到底应该放在哪里？

这里说的"调度"包含好几层意思：决定哪个计算块交给哪个 kernel、在哪个设备或队列上执行、以什么顺序提交、依赖关系怎么同步、遇到 if/while 的时候怎么走。传统方案里，这些逻辑大部分放在运行时的"通用调度引擎"里，现场读图、现场推导。IREE 的做法截然相反：把调度的结论尽量在编译期定死，写成一段字节码程序，运行时只负责按指令行事。我把这个区别想明白之后，才真正理解为什么 IREE 把产物做得这么小还能保持模型控制流完整。

1.1 传统部署方案把调度留在了运行时里

TFLite、ONNX Runtime 这类图解释器，模型文件本质是 FlatBuffer/Protobuf 格式的图描述，运行时拿到图之后按拓扑顺序逐节点执行。调度方案（比如算子融合、kernel 选择、多设备分配）很多是运行时根据 ops 类型做的，或者依赖外部 delegate。框架为了通用，运行时必须内置一大堆调度分支和兼容逻辑，体积天然小不了。TFLite Micro 为了上 MCU 做了裁剪，但本质上还是"通用解释器 + 图数据"，调度骨架依然住在运行时里。

这种架构有一个连带问题：产物不是自洽的。图描述是一份文件，kernel 可能是运行时内置库，也可能是单独动态库，调度策略则藏在运行时代码里。部署时任何一部分换版本，整个系统都要重新验证。我踩过的坑包括：图很小但连了一个很大的 kernel 库，或者运行时升级之后某个 op 的调度行为变了，线上模型结果对不上。

1.2 IREE 的立场：调度是编译问题，不是运行时问题

IREE 的核心思想是：模型在编译期被完整走一遍——哪些算子融合成同一个 dispatch、每个 dispatch 的 workgroup 规模、跑在哪个设备队列、依赖怎么串，都在这时候定下来。产物.vmfb里装的是已经排好顺序的 VM 字节码，字节码调用 HAL（硬件抽象层）接口把 kernel 发出去。运行时没有"图遍历"和"调度决策"的概念，只有一条一条执行的指令。

换句话说，"kernel A 必须等 kernel B 跑完"这种依赖，在产物里要么体现为指令先后顺序，要么体现为信号量 timepoint，而不是运行时现场分析出来的。调度逻辑烧进了二进制，执行实体（kernel 二进制）也烧进了同一个文件。这就是"把调度和执行一起烧进 10 KB 产物"这句话的真正含义：它不是一个营销口号，而是架构选择的结果。

2. 从输入图到 vmfb：五个阶段把调度和执行焊死

很多人一上来就看 IREE 的运行时，但其实精华在编译管线。整条链路是建立在 MLIR 之上的，大致可以切成五个阶段。以下按当前版本的主流流程来讲，具体 pass 名称在不同版本会略有出入，但骨架是稳定的。

2.1 前端与输入方言：MHLO/StableHLO/TOSA 只是原料

IREE 自己不直接吃 TensorFlow 或 PyTorch 的图，而是先转成 MLIR 里的标准张量方言。TF 走 iree-import-tf 转成 MHLO/StableHLO，PyTorch 走 torch-mlir 转换，TFLite 走 iree-import-tflite 转成 TOSA。这一步大多数人不会手动碰——官方提供了对应前端工具。但理解这一点很重要：IREE 之所以能接这么多框架，是因为它把"前端方言"和"后端设备"解耦了。中间的核心编译逻辑只认一份统一的输入形态。

2.2 flow 阶段：把计算切割成 dispatch region

这一步是调度的第一次成形。IREE 会分析整个计算图，把可以融合、可以切分的工作量划分成一个个 dispatch region。flow.dispatch.workgroups把一个 region 映射成一组 workgroup，每个 dispatch 最终对应一次 kernel 启动。比如一个add → mul → relu的链式结构，如果形状静态、布局匹配，就可以融合成一个 dispatch region，运行时只发一次 kernel，中间结果不出设备。

这个阶段很像编译器里的指令选择——它决定了"执行实体"的最小颗粒度。dispatch region 切得太碎，kernel 启动开销大；切得太粗，融合后的 kernel 可能寄存器压力爆掉或者分块不合理。IREE 的启发式会基于代价模型做选择，但代价模型不一定适合所有模型，所以我后来养成一个习惯：编完先跑一遍 benchmark，如果性能不对劲，再考虑用预处理 pass 干预融合。

2.3 stream 阶段：命令缓冲、时间点、队列同步

如果说 flow 决定了"什么能一起跑"，stream 就决定"以什么顺序跑、哪些可以重叠"。stream 方言引入stream.cmd.execute这样的命令块概念，里面装的是stream.cmd.dispatch、stream.cmd.fill、stream.cmd.copy这类命令，以及stream.timepoint这样的信号量描述。多设备场景下，不同队列之间通过 timepoint 建立依赖，允许 CPU 和 GPU 上的重叠执行在编译期就被安排好。

这个阶段其实就是"调度"的核心实现：命令缓冲的内容在编译期就录制好了。传统方案里，运行时要在提交前现场构建 command buffer，IREE 把这步提前到了编译期。运行时要做的事情就简化成：按顺序提交已录制好的命令块，在 timepoint 上做信号量等待。副作用是产物里天然包含了调度顺序，可读性差一点，但确定性和可复现性极强。

2.4 hal 与 vm 阶段：设备 ABI 和指令序列

hal 方言把 stream 里的调度动作映射到设备操作：hal.device.allocator拿缓冲、hal.buffer管理内存、hal.command_buffer提交命令、hal.executable对应实际 kernel 二进制。每个后端（llvm-cpu、Vulkan、CUDA、Metal）都要实现这层 ABI，负责把自己的 kernel 编译成目标格式。

真正把调度"焊死"的是 vm 阶段。IREE 把整个模型转换成一个带类型寄存器的字节码函数，所有 hal 调用变成vm.call.import，模型里的 while 循环变成字节码里的跳转指令，if/else 变成分支指令。换句话说，模型的运行时结构不再靠一张图来表达，而是被翻译成了一小段可以直接解释执行的指令序列。

我拿一个小模型做例子，它在 VM 里的大致形态是这样（不是精确的文本 IR，只是帮助理解）：

函数入口 分配输入/输出缓冲（hal.buffer.allocate） 创建命令缓冲（hal.command_buffer.create） 编码 dispatch（hal.command_buffer.dispatch） 提交命令缓冲（hal.command_buffer.submit + 信号量等待） 读回结果（hal.buffer.copy 或 mapping） 函数返回

如果模型里有循环，那么上面这段会被包在一个带条件跳转的循环体里，循环变量、累积状态都映射成 VM 寄存器。这就是为什么 IREE 对"模型内部有控制流"的场景特别友好——控制流不是 runtime 的通用能力，而是被编译成了产物的一部分。

2.5 产物内部结构：一个容器装了两样东西

最终.vmfb是一个 FlatBuffer 容器，里面至少包含三块内容：VM 字节码（调度编排逻辑）、嵌入的 HAL executable（实际 kernel 的机器码/SPIR-V/cubin），以及常量数据池。运行时加载 vmfb 时，先读字节码，再按需把 embedded executable 交给对应 HAL 驱动去准备（CPU 后端做重定位即可，GPU 后端把 SPIR-V 提交给驱动编译）。

这种"单文件自洽"的形态好处很直接：部署、签名、版本管理都围绕一个文件来做，不用再维护"图描述 + kernel 库 + 运行时版本"三方对齐。这也回答了标题里的问题——调度和执行确实被钉在同一个产物里了。

3. 10 KB 的量级怎么抠：字节码、Kernel 和裁剪三本账

标题说 10 KB 产物，我先把话说透：这指的是最小模型 + 静态形状 + 单后端 + 裁剪调试信息的理想下限，不是所有模型都能这么小。但理解这个下限怎么来的，对评估 IREE 特别有帮助。

3.1 一个 dispatch 在字节码里只值几条指令

IREE 的 VM 字节码设计得非常紧凑：操作码单字节、操作数按变长整型打包、寄存器编号是局部的。一次 dispatch 的开销，在字节码里只对应几条指令——分配缓冲、创建命令缓冲、编码 dispatch、提交、等待。对一个纯 elementwise 模型，整个字节码可能只有几十条指令，折合下来不到 2 KB。

调度逻辑是按"指令条数"计费，而不是按"运行时要实现的调度引擎"计费。传统解释器哪怕模型很小，运行时里也得带着一整套调度框架；IREE 把框架成本摊到编译期，产物的调度部分就只是具体问题的指令序列。这就是为什么产物能做小的根本原因。

3.2 Kernel 侧的体积：ELF vs SPIR-V

嵌入的 kernel 二进制是产物里的另一大块。llvm-cpu 后端把每个 dispatch 编译成一个函数，打进一个位置无关的 ELF 文件。一个f32+x这种 elementwise kernel 生成的机器码非常小，加上函数表、符号和元数据，可能就 5 KB 上下。如果是 ARM Thumb 指令集，还能更小。这里有个经验：同一份模型，llvm-cpu 内核通常比 SPIR-V 小，因为 SPIR-V 带完整类型描述和大量注解；CUDA cubin 也不小。所以"10 KB 产物"基本是 CPU 后端的叙事。

不同后端的单个 kernel 体积大概是这个量级（实测会因模型和参数浮动，仅供参考）：

3.3 让产物变小的编译参数清单

体积控制是有明确手段的，我整理了一张表。注意不同版本参数名可能略有出入，最好用iree-compile --help对一下当前版本：

真正跑起来的时候，我会先拿最朴素配置编译一遍，看ls -lh的体积，再逐步加功能，每加一个特性就量一次体积涨幅。这样能很清楚知道是哪项把产物撑大的。比如动态形状，代价是 host 端要生成一堆 shape 推导代码进字节码，模型不大体积却很容易冲过几十 KB。

4. 实操复现：用 iree-compile 把调度和执行装进一个文件

光讲架构容易飘，还是走一遍完整操作比较踏实。以下我用的版本对应官方最新稳定发布，命令在文档里都能找到对应说明。

4.1 环境准备与一个最小的 MLIR 模型

装两个包就够了：

pip install iree-compiler iree-runtime

然后写一个最小的 StableHLO 模型文件add.mlir：

func.func @add(%a: tensor<4xf32>, %b: tensor<4xf32>) -> tensor<4xf32> { %0 = stablehlo.add %a, %b : tensor<4xf32> return %0 }

真实模型一般不需要手写 MLIR，TensorFlow 和 PyTorch 都有对应的导入工具。这里用纯 MLIR 是为了把全链路看明白，不被前端转换的复杂度干扰。

4.2 编译、跑测、看大小

编译命令如下：

iree-compile add.mlir \ --iree-hal-target-backends=llvm-cpu \ --iree-llvmcpu-target-cpu-features=host \ --iree-hal-executable-debug-level=0 \ -o add.vmfb ls -lh add.vmfb

在我机器上，这个文件的体积大概在 10 KB 上下。就是标题说的那个量级。你可以拿这个文件继续跑一下：

iree-run-module --module=add.vmfb \ --device=local-task \ --function=add \ --input=4xf32=1,2,3,4 \ --input=4xf32=5,6,7,8

输出会是[6, 8, 10, 12]。这一步跑通，说明调度（指令顺序）和执行（内部 kernel）已经都进到单文件里，并且运行时的 local-task 驱动能正确解释执行。--device=local-task是 CPU 上最常用的驱动，底层用 IREE 自己的并行任务系统做 work-stealing 调度。

4.3 Python 运行时里加载 vmfb 调用

如果要在自己的程序里集成，用 Python 也很直接：

import iree.runtime as rt import numpy as np config = rt.Config(driver_name="local-task") ctx = rt.SystemContext(config=config) vm_module = rt.VmModule.from_flatbuffer_file("add.vmfb") ctx.add_module(vm_module) # 模块名默认一般为 module，函数名是 add；如果 dump 过可以看到真实名字 func = ctx.modules.module.add a = np.array([1.0, 2.0, 3.0, 4.0], dtype=np.float32) b = np.array([5.0, 6.0, 7.0, 8.0], dtype=np.float32) print(func(a, b))

不同小版本的 Python API 偶尔有微调，比如Config的构造方式、VmModule.from_flatbuffer_file的入参。遇到报错通常是因为你手头版本和示例不一致，查一下对应版本的 runtime 文档就好。

4.4 换 GPU 后端与交叉编译的快速体验

同一份add.mlir，换 Vulkan 后端编一次：

iree-compile add.mlir \ --iree-hal-target-backends=vulkan-spirv \ --iree-hal-executable-debug-level=0 \ -o add_vulkan.vmfb

对比一下体积，能明显感受到 SPIR-V kernel 比 ELF kernel 胖一圈。这就是我之前说的后端决定产物体积上限。

交叉编译到 ARM/RISC-V 的 MCU 时，llvm-cpu 后端支持指定目标 triple 和 CPU features，把产物编成目标架构的 ELF，再通过xxd -i add.vmfb > add_vmfb.c之类的方式转成 C 数组嵌进固件。运行时侧用裁剪过的 VM 解释器和最小 HAL 驱动加载这份字节码。这一步的坑主要在于 runtime 也要按目标平台交叉编译，建议直接参考官方 embedded 相关示例起步。

4.5 把 vmfb 反编译回指令看调度

我特别喜欢做的一件事：把编译好的 vmfb 反编译成文本，直接看调度指令。工具链里提供模块转储功能（我版本里类似iree-dump-module的工具），把字节码 dump 出来后，你能看到完整函数体、vm.call.import指向的 HAL 调用、循环跳转的 label。对一个带 while 的模型，你会清楚地看到字节码里有条件跳转指令——也就是说模型的控制流确实被编译成了产物的一部分，运行时不需要任何图解释逻辑。

这一步对排查奇怪行为非常有帮助。比如怀疑模型里有没被融合的 dispatch，dump 一下就知道到底有多少次 kernel 提交。

5. 我用 IREE 踩过的一串坑（以及绕坑姿势）

IREE 优点明显，坑也不少。下面这几个是我实际踩过、并且有代表性值得展开的。

5.1 vmfb 是后端专属产物，别幻想通用

.vmfb这个名字容易让人误以为是平台无关的字节码文件。字节码部分确实可以通用，但里面嵌入的 HAL executable 是严格绑定后端的：llvm-cpu 的 ELF 给 Vulkan 后端用不了一点。所以 CI 或构建脚本里，一定要按照目标设备分别编译产物，不能一份 vmfb 到处分发。我自己是建了一个"每设备一产物"的矩阵构建，把后端、优化参数、产物哈希一起打进版本号。

5.2 版本不匹配：VM ABI 的隐形锁

IREE 的 VM ABI 有明确的版本兼容窗口。编译器版本太新、运行时版本太旧，加载 vmfb 时会直接报 ABI 版本不匹配。这种问题在 Python 环境里最隐蔽，因为iree-compiler和iree-runtime是两个独立包，很容易升级其中一个另一个没跟上。建议一直保持两个包同版本安装，生产环境锁版本号。

5.3 动态形状会让体积和性能一起翻车

我一开始图省事，接口全用动态形状，想着灵活性高。结果产物体积直接从十几 KB 跳到几十上百 KB，吞吐还降了。原因是动态形状时，编译期无法把 shape 全部折叠进字节码，host 端必须生成大量 shape 推导、维度检查代码，同时 kernel 也要写通用的动态分支逻辑。后来我改成"大部分静态 + 必要时限幅（bound）"的策略，体积和性能都回来了。能静态尽量静态，这是 IREE 使用里最快见效的一条原则。

5.4 小模型跨设备调度可能倒亏性能

IREE 的多设备调度能力很强，可以一个模型内 CPU/GPU 协同、队列间用信号量同步。但能力强不等于任何模型都该用。我试过一个很小的 elementwise 模型，强行让部分 dispatch 走 GPU，结果设备间同步和拷贝开销比计算本身还大，整体比纯 CPU 慢一倍不止。经验法则：单算子都能在 CPU 上几微妙内算完的小模型，老老实实local-task单设备；只有模型足够大、并行度足够高，跨设备调度才划算。

5.5 调试工具链的几句经验

IREE 的调试手段比传统框架多一点，但需要一点耐心。排查编译问题我一般用--mlir-print-ir-after-all导出一大段中间结果，再配合grep找关键 pass 的输出；排查调度问题，看 stream 阶段 dump 的命令块结构；测性能用自带的iree-benchmark-module，并且一定要先 warmup 再取稳态数据，不然第一次调用里的 kernel 编译/缓存开销会被算进去，特别误导。

6. 什么样的情况适合押注 IREE，什么样的场景别硬上

作为选型参考，我把常见路线和 IREE 放在一起做个边界分析。这个表是我自己总结的，主要看"调度逻辑住在哪、运行时多大、控制流支持怎么样、学习成本多高"这几个维度：

我的选择经验分两种情况。如果你的目标是 MCU 类设备、闪存按 KB 算、模型里还带 while/if 这类动态控制流，IREE 的"编译期调度 + 迷你 VM"几乎是目前最匹配的路线——它把控制流和调度都变成了产物的一部分，运行时不用背解释器的包袱。如果项目周期很紧、团队没有 MLIR/编译器背景，只在单一平台上快速上线，那就别硬上 IREE，TF Lite / ONNX Runtime 的生态和资料会帮你省大量时间。

最后分享一个我现在固定的工作流程：拿到模型先转成静态形状，用 llvm-cpu 单后端编译一版 baseline，量体积、跑iree-benchmark-module记录时延；确认 baseline 没问题之后，再根据目标设备考虑换 GPU 后端、开微内核或者做多设备调度。vmfb 我始终当成烧录固件管理——每次发布都记录编译器版本、后端、优化参数和产物哈希。这样就算哪天线上出问题，也能很快定位是产物问题还是运行时版本问题。IREE 这套"把调度和执行一起烧进产物"的思路，实际用起来比表面看上去更值钱，前提是你愿意先吃透编译链路上这几个阶段。