note_AI_compiler/iree/02-architecture.md

阶段 2：IREE 架构总览与完整管线 ✅

2026-08-14。用 --mlir-print-ir-after-all 追踪一个 linalg matmul 从输入到 VM 的完整变换， 把 231 个 pass 划分成 7 个阶段，建立 IREE 整体的心智模型。

1. 一句话总结

IREE 是一个**「图 → 异步调度 → 设备代码 → 字节码」的完整编译栈**，每一层都用 MLIR dialect 表达： 输入方言（linalg/tosa/stablehlo）→ Flow → Stream → HAL → VM，其中 codegen（可执行体 lowering） 在 Stream 和 HAL 之间发生。

2. 完整 pipeline 全景（231 个 pass → 7 个阶段）

阶段	关键 dialect	代表性 pass	职责	TVM 对应
A 输入	(输入方言) + util + hal.tensor	WrapEntryPointsPass、InlinerPass	统一入口、包 ABI（buffer_view）	前端导入
B 设备	hal.device.target	AssignLegacyTargetDevicesPass、MaterializeTargetDevicesPass	目标设备分配	device plan
C Flow	flow	FormDispatchRegionsPass、OutlineDispatchRegionsPass	图级优化、fusion、dispatch 划分	Relay 图优化 + fusion
D Stream	stream	ConvertToStreamPass、ScheduleExecutionPass	异步调度（timepoint/fence）	（无，IREE 特色）
E Codegen	linalg/vector/llvm	TileAndFuse、GenericVectorization、Bufferize、ConvertToLLVM	可执行体 lowering 到设备代码	TE/TIR + schedule + codegen
F HAL	hal	ConvertToHALPass、SerializeAllExecutablesPass	设备抽象、可执行体序列化	runtime device 抽象
G VM	vm	（尾部 pass + 序列化）	字节码序列化	VM

完整 231 个 pass 的去重序列见 examples/matmul_pipeline_dump.txt （32702 行完整 dump，可用 grep 'IR Dump After' 提取 pass 序列）。

3. 关键 IR 变换（每个 dialect 边界）

边界 1：输入 → ABI（WrapEntryPointsPass）

用户写的 func(tensor) 被包成对外 util.func public，参数变成 !hal.buffer_view， 内部用 hal.tensor.import / hal.tensor.export 在 buffer_view 和 tensor 之间转换：

util.func public @matmul(%arg0: !hal.buffer_view, %arg1: !hal.buffer_view) -> !hal.buffer_view {
  %0 = hal.tensor.import %arg0 "input0" : !hal.buffer_view -> tensor<2x4xf32>
  %1 = hal.tensor.import %arg1 "input1" : !hal.buffer_view -> tensor<4x2xf32>
  %2 = util.call @_matmul(%0, %1) : (...) -> tensor<2x2xf32>
  %3 = hal.tensor.export %2 "output0" : tensor<2x2xf32> -> !hal.buffer_view
  util.return %3 : !hal.buffer_view
}

边界 2：Flow dispatch 划分（FormDispatchRegionsPass → OutlineDispatchRegionsPass）

linalg.matmul 被包进 flow.dispatch.region（→ 相当于 TVM 里划分 kernel）， 再 outline 成独立的 flow.executable，主函数里只剩一个 flow.dispatch 调用：

// FormDispatchRegionsPass 之后：
%6 = flow.dispatch.region -> (tensor<2x2xf32>) {
  %9 = linalg.matmul ins(%1, %3 : ...) outs(%8 : ...) -> tensor<2x2xf32>
  flow.return %9 : tensor<2x2xf32>
}

// OutlineDispatchRegionsPass 之后：
flow.executable private @matmul_dispatch_0 {
  flow.executable.export public @matmul_dispatch_0_matmul_2x2x4_f32 workgroups() -> (index, index, index) { ... }
  builtin.module {
    func.func @matmul_dispatch_0_matmul_2x2x4_f32(%arg0: !iree_tensor_ext.dispatch.tensor<...>, ...) {
      // 可执行体内部：linalg 计算 + dispatch.tensor.load/store
    }
  }
}
util.func public @matmul(...) {
  %2 = flow.dispatch @matmul_dispatch_0::@matmul_dispatch_0_matmul_2x2x4_f32(%0, %1) : (...) -> tensor<2x2xf32>
}

边界 3：Flow → Stream（ConvertToStreamPass）

可执行体变 stream.executable，绑定变 !stream.binding；主函数变成异步资源流：

// 可执行体函数签名：!stream.binding（抽象设备缓冲区）
func.func @matmul_dispatch_0_matmul_2x2x4_f32(%arg0: !stream.binding, %arg1: !stream.binding, %arg2: !stream.binding)

// 主函数：tensor → stream.resource + 设备亲和性标注
%1 = stream.tensor.import on(#hal.device.affinity<@__device_0>) %arg0 : !hal.buffer_view -> tensor<2x4xf32> in !stream.resource<external>{%0}
%7 = stream.tensor.dispatch on(#hal.device.affinity<@__device_0>) @matmul_dispatch_0::@matmul_dispatch_0_matmul_2x2x4_f32(%2, %5) : (...) -> tensor<2x2xf32> in !stream.resource<*>{%6}

边界 4：Stream 异步调度（ScheduleExecutionPass）

dispatch 被包进 stream.async.execute，返回 !stream.timepoint，再 await：

%results, %result_timepoint = stream.async.execute on(...) with(%0 as %arg2, %1 as %arg3) -> !stream.resource<...> {
  %4 = stream.async.dispatch @matmul_dispatch_0::@matmul_dispatch_0_matmul_2x2x4_f32(%arg2[...], %arg3[...]) : ...
  stream.yield %4 : !stream.resource<...>
} => !stream.timepoint
%2 = stream.timepoint.await %result_timepoint => %results : !stream.resource<...>

timepoint 是 IREE 异步模型的核心：类似 future/fence，表示「某批异步工作何时完成」， 用来表达数据依赖和跨设备同步。这是 TVM 没有的层。

边界 5：Codegen（可执行体 lowering，ConvertToLLVMPass）

可执行体里的 linalg matmul 被 tiling → vectorization → bufferize → LLVM：

// 最终：显式循环 + fmuladd + 向量 load
llvm.func @matmul_dispatch_0_matmul_2x2x4_f32(%arg0: !llvm.ptr, %arg1: !llvm.ptr, %arg2: !llvm.ptr) -> i32 {
  // 从 iree_hal_executable_dispatch_state_v0_t 取 bindings
  // ^bb1/^bb2/^bb3 显式循环，llvm.intr.fmuladd 做 MAC，vector<1xf32> 向量化
}

注意：codegen 在 Stream 和 HAL 之间发生（pass 130~199），因为可执行体要先编译成设备代码， 才能在 HAL 层被序列化。关键 pass：LLVMCPUSelectLoweringStrategyPass（选 tiling/向量化策略）、 TileAndDistributeToWorkgroups、LLVMCPUTileAndFuse、GenericVectorization、 IREEComprehensiveBufferize、ConvertToLLVM。

边界 6：Stream → HAL（ConvertToHALPass）

可执行体包成 hal.executable.variant，带 pipeline layout（缓冲区绑定）和 ordinal：

hal.executable private @matmul_dispatch_0 {
  hal.executable.variant public @embedded_elf_x86_64 target(#executable_target_embedded_elf_x86_64) {
    hal.executable.export public @matmul_dispatch_0_matmul_2x2x4_f32 ordinal(0) layout(#pipeline_layout)
      attributes {workgroup_size = [1 : index, 1 : index, 1 : index]}
    builtin.module { llvm.func @matmul_dispatch_0_matmul_2x2x4_f32(...) { ... } }
  }
}
// #pipeline_layout 定义了 3 个 storage_buffer 绑定（2 ReadOnly + 1 输出）

边界 7：最终 VM（字节码模块）

序列化后的 VM 模块：声明类型、依赖 hal 模块、import 一堆 hal.* 函数：

Required Types: i32, i64, !hal.allocator, !hal.buffer, !hal.buffer_view,
                !hal.command_buffer, !hal.device, !hal.executable, !hal.fence, !vm.buffer
Imported Functions:
  [ 7] hal.command_buffer.dispatch(...) -> ()
  [11] hal.device.queue.execute.indirect(...) -> ()
  [16] hal.fence.await(i32, i64, tuple<!vm.ref<?>>...) -> (i32)
  ...

运行时 iree-runtime 加载这个 VM 模块、解析 hal 模块 import、用 HAL driver 执行。

4. TVM 0.6 对照总结

IREE 层	作用	TVM 0.6 对应	差异点
Flow	图级优化 + dispatch 划分	Relay（fusion、算子划分）	划分出的 dispatch = 一个 kernel
Stream	异步调度	无	timepoint/fence 显式建模异步
Codegen	可执行体 lowering	TE/TIR + schedule + codegen	MLIR dialect lowering，非 schedule 原语
HAL	设备抽象	runtime device + codegen 的 device 层	更抽象，统一 CPU/GPU/Vulkan
VM	字节码	graph runtime / 早期 VM	紧凑字节码 + 模块依赖

5. 复现命令

# 完整 dump（含每个 pass 的 IR）
iree-compile --iree-hal-target-backends=llvm-cpu --mlir-print-ir-after-all \
  examples/matmul.mlir -o /tmp/matmul.vmfb 2> dump.txt

# 只看 pass 序列
grep -oE 'IR Dump After .*//' dump.txt | sed 's/IR Dump After //; s/ (.*//; s/ \/\/$//' \
  | awk '!seen[$0]++' | nl

# 只看某个 pass 之后的 IR
iree-compile --iree-hal-target-backends=llvm-cpu \
  --mlir-print-ir-after=ConvertToStreamPass --mlir-print-ir-module-scope \
  examples/matmul.mlir -o /tmp/matmul.vmfb 2>&1 | head -80

# 看最终 VM 模块
iree-dump-module /tmp/matmul.vmfb

6. 下一步

• 阶段 1（MLIR 基础）可选补：现在读懂了 dialect 转换，回头看 Operation/Value/Pass/PatternRewriting 会很快，可先快速过一遍再深入各层。
• 阶段 3（前端）、4（Flow+Stream）、5（HAL）、6（codegen）、7（运行时）逐层深入， 每一层都对着 src/iree/compiler/src/iree/compiler/Dialect/<X>/ 的源码读。