first commit

2026-08-20 17:20:09 +08:00
commit 8049b6f58d
14 changed files with 1649 additions and 0 deletions
@@ -0,0 +1,12 @@
+# 编译器学习文档
+
+> 个人编译器学习笔记汇总。每个编译器一个子目录，便于扩展到多个编译器。
+
+## 编译器目录
+
+- [iree/](iree/) —— IREE（构建于 MLIR 之上的编译器 + 运行时）
+
+## 当前进展
+
+- 正在学：IREE，规划见 [iree/00-plan.md](iree/00-plan.md)
+- 背景：熟悉 TVM 0.6（Relay / TE-TIR / schedule / codegen / runtime），曾接触过 IREE 的 codegen 部分
@@ -0,0 +1,120 @@
+# IREE 学习规划
+
+> 目标：系统掌握 IREE **整体**，建立完整的 pipeline 心智模型。
+> codegen 只是其中一个环节（恰好有经验，可快速过），不作为重点。
+
+## 0. 先回答：是否应该先学 MLIR？
+
+**结论：要学，但用「MLIR 基础 + 边拆 IREE 边补 MLIR」的方式，不要从头啃完 MLIR 教程。**
+
+理由：
+- IREE 的**每一层都是 MLIR dialect**（Flow、Stream、HAL、以及 codegen 用到的 Linalg/Vector），
+  不只是 codegen 依赖 MLIR。没有 MLIR 基础，整个 pipeline 都看不懂。
+- 但已有 TVM 0.6 经验，IR/pass/调度/访存等**编译器核心概念已具备**，MLIR 只是换表达方式，
+  不需要从零学编译原理，只需补「MLIR 的表达方式 + 基础设施」。
+
+**MLIR 需要掌握的核心（聚焦这些即可，不必面面俱到）：**
+- Operation / Value / Block / Region / SSA —— 对应熟悉的 IR node / expression
+- Dialect 与 OpBuilder
+- **Pass + Pattern Rewriting / Dialect Conversion** —— MLIR 的两种 pass 机制
+- 核心 dialect：`linalg` `tensor` `vector` `memref` `arith` `scf` `func`
+
+## 1. IREE 完整 pipeline（心智模型，先建立这个）
+
+```
+PyTorch / TF / JAX / ONNX
+      ↓  (torch-mlir / iree-import 系列工具)
+StableHLO / TOSA        ← MLIR 前端方言（跨框架的统一入口）
+      ↓
+IREE Flow dialect       ← 图级优化：type 推导、算子融合、划分 dispatch region（executable）
+      ↓
+IREE Stream dialect     ← 异步执行模型：timepoint / fence 调度，把图变成可异步执行的流
+      ↓
+IREE HAL dialect        ← 设备抽象：device / command buffer / allocator / executable 加载
+      ↓
+Codegen（executable）   ← Linalg/Vector/Tensor → LLVM IR / SPIR-V（设备相关代码生成）
+      ↓
+IREE VM                 ← 序列化成紧凑字节码，供运行时加载
+      ↓
+iree-runtime（C）       ← 加载 VM 模块 + HAL driver（CPU/CUDA/Vulkan/...）执行
+```
+
+## 2. 关键概念映射（TVM 0.6 → IREE/MLIR）
+
+| TVM 0.6 | IREE / MLIR | 说明 |
+|---|---|---|
+| Relay（图级 IR、fusion、算子融合） | Flow dialect + StableHLO/TOSA | 图级优化、形成 dispatch region |
+| TE / TIR + schedule 原语（split/fuse/reorder/unroll） | Linalg + Vector + tiling/vectorization | tensor-level → loop-level 代码生成 |
+| codegen（LLVM/CUDA/OpenCL） | 翻译到 LLVM / SPIR-V | 后端代码生成 |
+| runtime（graph runtime / 早期 VM） | HAL + VM | 设备抽象 + 字节码执行 |
+| （TVM 无直接对应） | Stream dialect | 异步调度，IREE 特色之一 |
+
+## 3. 学习阶段（系统覆盖整体）
+
+### 阶段 0：环境准备
+- 安装 IREE（推荐 nightly Python 包 `iree-compiler` + `iree-runtime`，或从源码编译）
+- 跑通一个最小例子：`iree-compile` 编译 → `iree-run-module` 执行
+- 熟悉工具链：`iree-compile` / `iree-run-module` / `iree-dump-module`
+
+### 阶段 1：MLIR 基础（IREE 全流程的地基）
+- MLIR 核心概念：Operation / Value / Block / Region / Dialect / Type / Attribute
+- 核心 dialect：`linalg` `tensor` `vector` `memref` `arith` `scf` `func`
+- Pass 与 Pattern Rewriting / Dialect Conversion
+- 工具：`mlir-opt` / `mlir-translate`，用 `--mlir-print-ir-before-all` / `--mlir-print-ir-after-all` 观察 pass 效果
+- 参考：MLIR 官方 Toy tutorial + 语言参考 + linalg 文档
+
+### 阶段 2：IREE 架构总览与完整管线
+- 用 `--mlir-print-ir-after-all` 追踪一个简单模型从输入到 VM 的完整变换
+- 认识每一层 dialect 的职责与边界（对上「心智模型」图）
+- 理解各层之间如何衔接（输入 → Flow → Stream → HAL → executable → VM）
+
+### 阶段 3：前端输入（StableHLO / TOSA）
+- 各框架如何进入 IREE：torch-mlir、iree-import-tf / iree-import-onnx
+- StableHLO 与 TOSA 的角色、差异、如何选择
+- 输入的图结构 → 进入 Flow 前的形态
+
+### 阶段 4：编译中段（Flow + Stream）
+- Flow：type 推导、fusion、dispatch region（executable）的形成
+- Stream：异步执行模型，timepoint / fence 如何表示数据依赖与调度
+- 图级优化与 TVM Relay 的对照
+
+### 阶段 5：HAL 设备抽象
+- device / command buffer / allocator / executable 的抽象设计
+- HAL 如何屏蔽 CPU / CUDA / Vulkan 等后端差异
+- 与 TVM codegen 的 device 抽象对照
+
+### 阶段 6：Codegen（可执行体 lowering，有经验可快进）
+- executable 的 lowering：Linalg → Vector → LLVM IR / SPIR-V
+- 后端选择：LLVMCPU、LLVMGPU(CUDA)、SPIR-V(Vulkan) 等
+- tile-and-fuse / tiling / vectorization / 访存优化
+- 与 TVM TE/TIR + schedule 的对应关系（你已熟，重点是补 MLIR 侧的机制）
+
+### 阶段 7：运行时（VM + iree-runtime）
+- IREE VM 字节码格式与执行模型
+- iree-runtime（C）：模块加载、HAL driver、调度执行
+- 与 TVM runtime 对比
+
+### 阶段 8：实践项目（巩固，可选）
+- 写一个自定义 pass（例如自定义 fusion / tiling pass）
+- 加一个自定义 backend，或写一个自定义 dialect 练手
+- 追踪并优化一个具体算子的端到端 lowering
+
+## 4. 学习资源
+
+- **IREE 官方文档**：<https://iree.dev/>
+- **IREE 代码库**：<https://github.com/iree-org/iree>
+  - 中段：`compiler/src/iree/compiler/Dialect/Flow/`、`.../Dialect/Stream/`、`.../Dialect/HAL/`
+  - codegen：`compiler/src/iree/compiler/Codegen/`
+- **MLIR 官方文档**：<https://mlir.llvm.org/>
+- **MLIR Toy tutorial**：<https://mlir.llvm.org/docs/Tutorials/Toy/>
+- **Linalg 文档**：<https://mlir.llvm.org/docs/Dialects/Linalg/>
+- **论文**：*Composable and Modular Code Generation in MLIR*（IREE codegen 的核心设计）
+- **torch-mlir**：<https://github.com/llvm/torch-mlir>（前端 → StableHLO/TOSA 参考）
+
+## 5. 学习方法建议
+
+- **以 IR dump 为主线**：每个阶段都用 `--mlir-print-ir-after-all` / `--mlir-print-ir-before-all`
+  看 IR 变化，这是理解 pass pipeline 最快的方式。
+- **对照 TVM**：每遇到一个 IREE 概念，都问自己「这在 TVM 里对应什么」，把已有经验迁移过来。
+- **按层推进，不跳层**：前端 → 中段 → HAL → codegen → 运行时，每层弄懂职责和边界后再进下一层。
+  codegen 有经验，可以快速过、只补 MLIR 侧的机制，但不要因为它熟就跳过其他层。
@@ -0,0 +1,107 @@
+# 阶段 0：环境准备 ✅
+
+> 2026-08-13 完成。记录 IREE 环境搭建过程，方便复现。
+
+## 1. 安装方式
+
+- **venv**（不用 conda，见下），位置：`.venv/`（项目根目录）
+- **PyPI 镜像**：清华源（pypi.org 实测仅 ~17-82KB/s，太慢；清华源 ~1.1MB/s）
+  - 配置写在 `.venv/pip.conf`：`index-url = https://pypi.tuna.tsinghua.edu.cn/simple`
+
+## 2. 安装的包
+
+```bash
+python3 -m venv .venv
+.venv/bin/pip install --upgrade pip
+.venv/bin/pip install iree-base-compiler iree-base-runtime
+```
+
+| 包 | 版本 | 说明 |
+|---|---|---|
+| `iree-base-compiler` | 3.11.0 | 编译器（含 `iree-compile`、`iree-opt` 等） |
+| `iree-base-runtime` | 3.11.0 | 运行时（含 `iree-run-module`、`iree-dump-module` 等） |
+
+### ⚠️ 包名要点（重要，坑）
+
+IREE 在 2024-11 前后重构了 Python 包：
+- 老的 `iree-compiler` / `iree-runtime`（日期版本号）已**停在 20241104.1068**，不再更新
+- 新的用**语义化版本号**，包名是 `iree-base-compiler` / `iree-base-runtime`（当前 3.11.0）
+- 没有 `iree-tools` 这个包，CLI 工具直接随 base 包分发
+
+## 3. 工具链
+
+`iree-compile` / `iree-run-module` / `iree-dump-module` / `iree-opt` / `iree-import-onnx` /
+`iree-benchmark-module` 等，都在 `.venv/bin/` 下。
+
+- 编译器版本：`IREE compiler 3.11.0rc20260316`，LLVM 23.0.0git
+- 注意：`iree-run-module` 是 C 工具，**不吃 `--version`**（报 `flag not recognized`），这是正常的
+
+## 4. 最小端到端验证
+
+文件：`examples/matmul.mlir`（linalg 矩阵乘 2x4 @ 4x2 → 2x2）
+
+```bash
+# 编译
+iree-compile --iree-hal-target-backends=llvm-cpu examples/matmul.mlir -o examples/matmul.vmfb
+# 运行
+iree-run-module --module=examples/matmul.vmfb --function=matmul \
+  --input=2x4xf32=1,2,3,4,5,6,7,8 --input=4x2xf32=1,2,3,4,5,6,7,8
+# 输出
+2x2xf32=[50 60][114 140]   # 与手算一致 ✓
+```
+
+### 编译时的提示
+
+`iree-compile` 会提示「未指定 CPU 目标」，建议加 `--iree-llvmcpu-target-cpu=host`
+（针对本机 CPU 生成最优代码）或 `=generic`（通用、可移植）。学习阶段可用默认值，不影响正确性。
+
+## 5. 硬件 / 后续
+
+- CPU：16 核 x86_64
+- GPU：NVIDIA RTX 4070 Ti SUPER（驱动 580.173.02）
+  - **CUDA backend 尚未配置**：base runtime 默认只有 CPU（local-task）。CUDA/Vulkan 后端后续需要：
+    - 编译器侧：`iree-compile --iree-hal-target-backends=cuda`（base 包可能不含 CUDA 代码生成，需确认）
+    - 运行时侧：CUDA HAL driver（可能需源码编译或插件）
+  - 学习阶段先用 CPU 即可，阶段 6 深入 codegen 时再按需补 GPU 后端
+- 无 Vulkan（未装 vulkaninfo），后续要用 Vulkan 后端需先装 Vulkan 驱动/loader
+
+## 6. 常用命令速查
+
+```bash
+# 激活环境
+source .venv/bin/activate
+# 编译（CPU）
+iree-compile --iree-hal-target-backends=llvm-cpu in.mlir -o out.vmfb
+# 运行
+iree-run-module --module=out.vmfb --function=<fn> --input=<shape>xtype=vals
+# 看编译后的模块结构 / 反汇编
+iree-dump-module --module=out.vmfb
+# 看完整 lowering 管线（学习核心工具）
+iree-compile --mlir-print-ir-after-all in.mlir -o out.vmfb
+```
+
+## 7. 源码（读 vs 构建）
+
+区分三件事：
+
+| 目的 | 需要什么 | 状态 |
+|---|---|---|
+| 跑 IREE、看 pipeline | pip 二进制包 | ✅ 已装 |
+| 读 IREE 源码 | `git clone`（只读） | ✅ 已 clone |
+| 改代码 + 断点调试 | 源码构建 | ⬜ 暂缓（阶段 8 / 需要 CUDA 时） |
+
+**已 clone**（只读、浅克隆、不拉子模块）：
+
+```bash
+cd src && git clone --depth 1 --single-branch https://github.com/iree-org/iree.git
+# 位置: src/iree/（约 211MB）
+```
+
+关键目录：
+- `compiler/src/iree/compiler/Dialect/` —— Flow / Stream / HAL / VM / VMVX / LinalgExt / TensorExt / Encoding / Util
+- `compiler/src/iree/compiler/Codegen/` —— LLVMCPU / LLVMGPU / Common / Dialect
+- `runtime/src/iree/` —— hal / vm / task / async / base
+
+**从源码构建（暂缓，先了解成本）**：需 clang + ninja，几十 GB 磁盘、~1 小时编译；
+只 debug 编译器可只构建 `iree-compile` 目标。学习阶段用 `--mlir-print-ir-after-all` 看 pipeline 即可，
+不必构建。到阶段 8 实践（写自定义 pass）或要 CUDA 后端时再构建。
@@ -0,0 +1,174 @@
+# 阶段 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](../../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 之间转换：
+
+```mlir
+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` 调用：
+
+```mlir
+// 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`；主函数变成异步资源流：
+
+```mlir
+// 可执行体函数签名：!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`：
+
+```mlir
+%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：
+
+```mlir
+// 最终：显式循环 + 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：
+
+```mlir
+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. 复现命令
+
+```bash
+# 完整 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>/` 的源码读。
@@ -0,0 +1,215 @@
+# 阶段 1：MLIR 速查（对应规划阶段 1）
+
+> 2026-08-14。面向「已有 TVM 经验、看过 IREE pipeline」的 MLIR 速查。
+> 重点是 dialect 概念 + 各 dialect 职责，配合阶段 2 的 pipeline dump 一起看。
+
+## 1. MLIR 的核心思想
+
+MLIR 不是一个固定 IR，而是一个 **IR 框架**：它定义了一套通用的 IR 结构（Operation/Block/Region/Type），
+然后让你**自己定义 dialect**。每个 dialect 定义自己的 op、类型、属性。
+
+- **builtin**：最核心的部分（module、基础类型）
+- **dialect**：一组相关的 op/type/attr 的集合（如 `arith`、`linalg`、`flow`）
+- **op**：具体操作（如 `arith.addf`、`linalg.matmul`）
+
+IREE 的每一层都是一个 dialect：Flow、Stream、HAL、VM 都是自定义的 MLIR dialect。
+
+## 2. IR 结构核心概念
+
+| 概念 | 是什么 | 类比 TVM |
+|---|---|---|
+| **Operation (op)** | 基本单位：名字 + operands + results + attributes + regions | IR 节点 / expr |
+| **Value** | SSA 值，只能赋值一次；是 op 的 result 或 block 的 argument | 变量/临时值 |
+| **Block** | 顺序执行的一串 op，可带 block argument，以 terminator 结尾 | 基本块 |
+| **Region** | 一个或多个 block 的列表，支持嵌套（实现控制流） | 子图/作用域 |
+| **Type** | 值的类型：`i32` `f32` `tensor<2x4xf32>` `memref<...>` `vector<...>` | dtype |
+| **Attribute** | 编译期常量/元数据：`arith.constant 5 : i32` 里的 5，或 op 的 `{...}` | 常量/annotations |
+| **Symbol** | 命名实体（函数名、全局名），通过符号引用 | 函数名 |
+
+### Op 的通用语法（重点，读 IR 必会）
+
+```mlir
+%result = linalg.matmul ins(%lhs, %rhs : tensor<2x4xf32>, tensor<4x2xf32>)
+                        outs(%init : tensor<2x2xf32>) -> tensor<2x2xf32>
+```
+
+- `%result` = 结果 SSA 值；`-> tensor<2x2xf32>` = 结果类型
+- `ins` / `outs` = 输入 / 输出操作数（linalg 特有的命名，其他 op 直接用位置参数）
+- `(%lhs, %rhs : type1, type2)` = 操作数及其类型
+
+## 3. 什么是 Dialect？
+
+**Dialect = 一组相关的 op、类型、属性的集合**，定义在同一个命名空间下。例如：
+- `arith.addf` 是 `arith` dialect 里的 op，名字用 `arith.` 前缀
+- `linalg.matmul` 是 `linalg` dialect 里的 op
+- `flow.dispatch.region` 是 IREE 的 `flow` dialect 里的 op
+
+为什么要分 dialect？因为**不同抽象层次需要不同的"语言"**：
+- 高层用 `linalg`（表达"这是一个矩阵乘"的语义）
+- 低层用 `llvm`（表达"这是 fmuladd + 循环 + 指针"）
+- 图级用 `flow`、异步用 `stream`、设备用 `hal`
+
+编译过程就是**从高层 dialect 逐层 lowering 到低层 dialect**（这正是阶段 2 看到的 7 个阶段）。
+
+## 4. 基础 dialect（IR 的积木）
+
+### `func` —— 函数
+```mlir
+func.func @matmul(%arg0: tensor<2x4xf32>) -> tensor<2x2xf32> {
+  %0 = func.call @helper(%arg0) : (tensor<2x4xf32>) -> tensor<2x2xf32>
+  func.return %0 : tensor<2x2xf32>
+}
+```
+- 定义函数 `func.func`、调用 `func.call`、返回 `func.return`
+
+### `arith` —— 算术/比较/转换
+```mlir
+%0 = arith.constant 5 : i32
+%1 = arith.addf %a, %b : f32          // 浮点加
+%2 = arith.addi %x, %y : i32          // 整数加
+%3 = arith.cmpf ogt, %a, %b : f32     // 浮点比较
+%4 = arith.mulf %a, %b : f32          // 浮点乘
+```
+- 标量算术的"标准库"，几乎所有计算最终都落到这里
+
+### `scf` —— 结构化控制流（for / if / while）
+```mlir
+%0 = scf.for %i = %lb to %ub step %step iter_args(%acc = %init) -> (tensor<...>) {
+  %1 = arith.addi %i, %acc : i32
+  scf.yield %1 : i32
+}
+%2 = scf.if %cond -> (f32) {
+  scf.yield %a : f32
+} else {
+  scf.yield %b : f32
+}
+```
+- "结构化"= 带明确的循环语义，便于分析和变换（IREE tiling 产生的就是 `scf.for`）
+
+### `cf` —— 非结构化控制流（底层分支）
+```mlir
+cf.br ^bb1(%0 : i32)                  // 无条件跳转
+cf.cond_br %cond, ^bb1(%0 : i32), ^bb2  // 条件跳转
+```
+- 最低层的控制流，`scf.for` 最终会 lower 成 `cf.br` 循环
+
+### `affine` —— 仿射循环与访存
+```mlir
+%0 = affine.for %i = 0 to 10 {
+  %1 = affine.load %memref[%i, %i + 1] : memref<10x10xf32>
+}
+```
+- 循环边界和下标必须是仿射表达式（`i`、`i+1`、`2*i+3`），可做 polyhedral 分析
+- IREE 历史上有用，现在更多用 scf + linalg
+
+### `index` 类型
+- `index` 是"可移植的整数"，循环变量、下标、shape 都用它
+- 编译到具体目标时再决定是 i32 还是 i64
+
+## 5. 数据/计算 dialect（IREE codegen 的核心）
+
+### `tensor` —— 不可变张量（值语义）
+```mlir
+%0 = tensor.empty() : tensor<2x2xf32>              // 建空张量
+%1 = tensor.extract_slice %t[0, 0][2, 2][1, 1] : tensor<4x4xf32> to tensor<2x2xf32>
+%2 = tensor.insert_slice %s into %t[0, 0][2, 2][1, 1] : tensor<2x2xf32> into tensor<4x4xf32>
+%3 = tensor.extract %t[0, 0] : tensor<2x2xf32>     // 取标量
+```
+- **值语义**：不可变，像纯函数式。编译期的高层抽象
+
+### `memref` —— 内存缓冲区（可变的 buffer）
+```mlir
+%0 = memref.alloc() : memref<2x2xf32>
+%1 = memref.load %0[%i, %j] : memref<2x2xf32>       // 读
+memref.store %v, %0[%i, %j] : memref<2x2xf32>       // 写
+```
+- 可变，带 layout/stride，对应"真实内存"。**bufferize 后 tensor 变 memref**
+
+### `vector` —— 固定大小 SIMD 向量
+```mlir
+%0 = vector.load %memref[%i, %j] : memref<2x2xf32>, vector<4xf32>
+%1 = vector.transfer_read %t[0, 0], %pad : tensor<4x4xf32>, vector<4xf32>
+```
+- 固定大小的向量，用于表达 SIMD 指令（IREE 向量化后产生 `vector` op）
+
+### `linalg` —— 结构化计算（**IREE 的心跳**）
+```mlir
+// 命名 op：语义明确的高层计算
+%0 = linalg.matmul ins(%lhs, %rhs : tensor<2x4xf32>, tensor<4x2xf32>)
+                   outs(%init : tensor<2x2xf32>) -> tensor<2x2xf32>
+%1 = linalg.fill ins(%cst : f32) outs(%empty : tensor<2x2xf32>) -> tensor<2x2xf32>
+
+// generic op：自定义的"结构化计算"，用 region 定义元素计算 + 迭代空间
+%2 = linalg.generic {
+  indexing_maps = [...],
+  iterator_types = ["parallel", "parallel", "reduction"]
+} ins(%a : tensor<...>) outs(%b : tensor<...>) {
+^bb0(%in: f32, %out: f32):
+  %sum = arith.addf %in, %out : f32
+  linalg.yield %sum : f32
+}
+```
+- **关键思想**：linalg 把"计算"和"数据如何存/遍历"分离。
+  它只声明"这是矩阵乘"或"元素级计算 + 迭代空间"，具体 tiling/向量化由后续 pass 做。
+  这就是 IREE codegen 的起点（对应 TVM 的 TE 层，但声明式更强）。
+
+## 6. 易混淆概念对照（重点）
+
+| 对比 | 区别 |
+|---|---|
+| `tensor` vs `memref` | 值语义（不可变）vs 内存（可变）。编译期用 tensor，bufferize 后变 memref |
+| `scf.for` vs `affine.for` | 结构化（可标并行/归约）vs 仿射（可 polyhedral 分析） |
+| `arith` vs `math` | 基本算术（add/mul/cmp）vs 数学函数（`math.sqrt`、`math.exp`、`math.absf`） |
+| `tensor.empty` vs `memref.alloc` | 抽象空张量（值）vs 真实内存分配 |
+
+## 7. Pass 基础设施（看懂 pipeline 的关键）
+
+**Pass** = 遍历 IR 做变换的基本单位。IREE 那 731 次 pass 就是一个个变换。
+
+两种主要机制（深入各层源码必懂）：
+
+### (1) Pattern Rewriting（局部重写）
+声明式：写"匹配 pattern → 替换成新 op"的规则，引擎（如 `applyPatternsAndFoldGreedily`）自动遍历应用。
+适合局部的、等价的化简。
+
+### (2) Dialect Conversion（跨 dialect 转换）
+三件套：`ConversionTarget`（哪些 op 合法/非法）+ `TypeConverter`（类型怎么换）+
+patterns（非法 op 怎么重写成合法 op）。
+适合把整个 dialect 换成另一个（如 Flow→Stream、tensor→memref）。
+
+> 你在 pipeline 里看到的 `ConvertToStreamPass`、`ConvertToLLVMPass`、`IREEComprehensiveBufferizePass`
+> 都是 Dialect Conversion；`CanonicalizerPass` 是 Pattern Rewriting。
+
+### 常用观察工具
+
+```bash
+iree-compile --mlir-print-ir-after-all in.mlir -o out.vmfb 2> dump.txt   # 每个 pass 后 dump
+iree-compile --mlir-print-ir-after=ConvertToStreamPass --mlir-print-ir-module-scope in.mlir -o out.vmfb 2>&1  # 只看某个 pass
+```
+
+## 8. 与 TVM 0.6 对照
+
+| MLIR | TVM 0.6 |
+|---|---|
+| Operation | IR node / expr |
+| Pass | pass / IR visitor |
+| Dialect Conversion | 逐层 lowering（Relay→TE→TIR） |
+| Pattern Rewriting | pass 里的 rewrite 规则 |
+| linalg.matmul/generic | TE 的 compute + 描述 |
+| scf.for + vector | TIR 里的 loop + vectorize |
+| memref | 缓冲区/物理内存 |
+
+## 9. IREE 自己的 dialect（概览，详细见各层笔记）
+
+| dialect | 职责 | 关键 op（在 dump 里见过） |
+|---|---|---|
+| `util` | 通用工具（函数/全局包装） | `util.func`、`util.global`、`util.call`、`util.return` |
+| `flow` | 图级优化 + dispatch 划分 | `flow.dispatch.region`、`flow.dispatch`、`flow.executable` |
+| `stream` | 异步执行模型 | `stream.executable`、`stream.binding`、`stream.tensor.import`、`stream.async.execute`、`stream.timepoint.await` |
+| `hal` | 设备抽象 | `hal.executable`、`hal.buffer_view`、`hal.command_buffer`、`hal.tensor.import/export` |
+| `vm` | VM 字节码 | `vm.module`、`vm.func`、`vm.call` |
+| `iree_tensor_ext` | dispatch 张量 + 屏障 | `iree_tensor_ext.dispatch.tensor.load/store`、`iree_tensor_ext.compute_barrier.start/end` |
+| `iree_linalg_ext` | 扩展 linalg op（softmax/attention 等） | 复杂算子的结构化 op |
+
+> 这些 dialect 的源码在 `src/iree/compiler/src/iree/compiler/Dialect/<X>/`，
+> 每个 dialect 目录下有 `*.td`（op 定义，TableGen 格式）和 `*.cpp`（实现）。
@@ -0,0 +1,159 @@
+# 阶段 3：前端输入（StableHLO / TOSA）✅
+
+> 2026-08-15。理解模型如何进入 IREE，以及两个统一的输入方言如何收敛到 linalg。
+> 对应规划阶段 3。源码位置：`src/iree/compiler/plugins/input/`。
+
+## 1. 模型怎么进入 IREE
+
+IREE 不直接吃框架模型，而是通过**导入工具**先把框架模型转成 MLIR 的「输入方言」，再走统一的编译管线：
+
+| 框架 | 导入工具 | 产生的输入方言 |
+|---|---|---|
+| PyTorch | torch-mlir / iree-turbine | Torch → StableHLO 或 TOSA |
+| TensorFlow | iree-import-tf（XLA 导出） | StableHLO |
+| JAX | jax.export / XLA | StableHLO |
+| ONNX | iree-import-onnx | TOSA |
+| TFLite | iree-import-tflite | TOSA |
+
+> 当前环境只装了 `iree-import-onnx`（base 包自带）。TF 导入器需要额外的 `iree-tools-tf` 包，
+> torch 前端需要 `iree-turbine`（依赖 torch）。后面阶段 3 深入时按需补装。
+
+## 2. 两个统一的输入方言
+
+IREE 支持多个输入方言，但**两个主力是 StableHLO 和 TOSA**，它们各自代表一条血统：
+
+| | StableHLO | TOSA |
+|---|---|---|
+| 全称 | Stable High-Level Ops | Tensor Operator Set Architecture |
+| 血统 | XLA / HLO（TF、JAX、PyTorch-XLA） | Arm TOSA 规范（ONNX、TFLite） |
+| 风格 | 图级高层 op（dot/transpose/select/broadcast） | 精确 spec 的算子集（max_pool2d/…，带 kernel/stride/pad 属性） |
+| 特点 | 保留框架 `{name=...}` 属性，op 粒度粗 | 每个 op 有精确语义规范，量化友好 |
+
+**为什么要两个？** 因为它们分别服务于不同的框架生态（XLA 系 vs ONNX/TFLite 系），
+但**最终都收敛到同一个内部表示（linalg / linalg-ext）**，所以下游编译管线是共享的。
+
+## 3. IR 示例
+
+### StableHLO（fullyconnected.mlir，一个全连接层）
+```mlir
+func.func @main(%arg0: tensor<1x5xf32>, %arg1: tensor<1x5x3x1xf32>) -> tensor<5x1x5xf32> {
+  %0 = stablehlo.transpose %arg0, dims = [1, 0] {name = "transpose.41"} : ...
+  %3 = stablehlo.broadcast_in_dim %2, dims = [0, 1] {name = "broadcast.56"} : ...
+  %29 = stablehlo.dot %28, %cst_11, precision = [DEFAULT, DEFAULT] {name = "dot.62"} : ...
+  %31 = stablehlo.add %29, %30 {name = "add.65"} : ...
+  %32 = stablehlo.maximum %25, %31 {name = "maximum.68"} : ...
+  %34 = stablehlo.select %23, %24, %33 {name = "select.71"} : ...
+}
+```
+- 高层图级 op，保留 TF 原图的 `{name = "..."}` 属性
+
+### TOSA（max_pool.mlir）
+```mlir
+%result = tosa.max_pool2d %0 {kernel = array<i64: 2, 2>, stride = array<i64: 1, 1>,
+                              pad = array<i64: 0, 0, 0, 0>} : (tensor<1x2x4x1xi8>) -> tensor<1x1x3x1xi8>
+```
+- 算子 + 显式属性（kernel/stride/pad），语义精确
+
+## 4. 前端转换：输入方言 → linalg
+
+用 `--mlir-print-ir-after-all` 追踪 StableHLO 模型的转换，前端约 19 个 pass，
+其中**最关键的两个**：
+
+| pass | 作用 |
+|---|---|
+| `ConvertStableHloToLinalgExt` | StableHLO → LinalgExt（IREE 扩展的 linalg） |
+| `ConvertStableHloToIreeInputDialects` | StableHLO → IREE 输入方言（linalg 等） |
+
+转换前后对照（`stablehlo.dot` → `linalg.generic`）：
+
+```mlir
+// 之前：stablehlo（高层语义）
+%29 = stablehlo.dot %28, %cst_11, precision = [DEFAULT, DEFAULT] {name = "dot.62"} : ...
+
+// 之后：linalg.generic（结构化计算 + 显式迭代空间）
+%1 = linalg.generic {indexing_maps = [#map, #map1],
+                     iterator_types = ["parallel", "parallel", "parallel"]}
+     ins(%expanded : tensor<5x1xf32>) outs(%0 : tensor<5x1x5xf32>) attrs = {name = "broadcast.56"} {
+^bb0(%in: f32, %out: f32):
+  linalg.yield %in : f32
+} -> tensor<5x1x5xf32>
+
+// stablehlo.compare GT → linalg.generic + arith.cmpf
+^bb0(%in: f32, %out: i1):
+  %45 = arith.cmpf ogt, %in, %cst : f32
+  linalg.yield %45 : i1
+```
+
+**关键点**：`stablehlo.*` 是"命名高层 op"，`linalg.generic` 是"结构化计算"——用
+`indexing_maps`（仿射映射描述数据布局）+ `iterator_types`（parallel/reduction）+
+region（元素级计算）显式描述计算。这就是阶段 2 里看到的 codegen 起点的由来。
+
+## 5. 源码结构（plugins/input/）
+
+前端转换是**插件式**的（IREE 2024 重构后，输入转换拆成独立 plugin）：
+
+```
+src/iree/compiler/plugins/input/
+├── StableHLO/
+│   ├── Conversion/           # StableHLO → linalg/linalg-ext
+│   │   ├── StableHLOToIREEInputDialects.cpp   # 主转换
+│   │   ├── StableHLOToLinalgExt.cpp
+│   │   ├── LegalizeControlFlow.cpp / LegalizeShapeComputations.cpp
+│   │   └── VerifyCompilerInputLegality.cpp
+│   └── PluginRegistration.cpp
+├── TOSA/
+│   ├── InputConversion/
+│   │   ├── TosaToLinalgExt.cpp
+│   │   ├── StripSignedness.cpp / Converti48Toi64.cpp  # TOSA 特有预处理
+│   │   └── VerifyCompilerTOSAInputLegality.cpp
+│   └── PluginRegistration.cpp
+└── Torch/                    # torch-mlir（Torch → StableHLO/TOSA）
+```
+
+## 6. 完整输入 pass 序列（StableHLO 模型，去重后前 19 个）
+
+```
+CheckVHLOStableHloMixUsage → VhloToVersionPass → VhloLegalizeToStablehloPass   # VHLO 版本化 → StableHLO
+→ Canonicalizer → StableHLOCanonicalize → CSE                                # 化简
+→ LegalizeStableHLOCustomCalls → LegalizeControlFlow → FlattenTuplesInSCF    # 控制流/自定义调用
+→ StableHLOToStableHLOPreprocessing → ShapeToShapeLowering → ConvertShapeToStandard  # shape 计算
+→ Inliner → LegalizeShapeComputations
+→ ConvertStableHloToLinalgExt → ChloLegalizeToStablehloPass                   # ★ StableHLO → linalg-ext
+→ ConvertStableHloToIreeInputDialects                                         # ★ → IREE 输入方言
+→ ReconcileUnrealizedCastsPass → VerifyCompilerStableHloInputLegality         # 校验
+```
+
+> 之后（pass 20+）就接上阶段 2 的 IREEImportPublicPass → WrapEntryPointsPass → … 主管线。
+
+## 7. 与 TVM 0.6 对照
+
+| IREE 前端 | TVM 0.6 |
+|---|---|
+| iree-import-tf / torch-mlir / iree-import-onnx | 各框架的 relay importer |
+| StableHLO / TOSA | Relay IR（图级） |
+| ConvertStableHloToIreeInputDialects（→linalg） | Relay → TE 的 lowering |
+| linalg.generic | TE 的 compute |
+
+## 8. 实操命令
+
+```bash
+# 编译 StableHLO 模型，看完整前端转换
+iree-compile --iree-input-type=stablehlo --iree-hal-target-backends=llvm-cpu \
+  --mlir-print-ir-after-all src/iree/tests/e2e/stablehlo_models/fullyconnected.mlir -o /tmp/fc.vmfb 2> dump.txt
+
+# 只看 StableHLO → linalg 的关键转换
+iree-compile --iree-input-type=stablehlo --iree-hal-target-backends=llvm-cpu \
+  --mlir-print-ir-after=ConvertStableHloToIreeInputDialects --mlir-print-ir-module-scope \
+  src/iree/tests/e2e/stablehlo_models/fullyconnected.mlir -o /tmp/fc.vmfb 2>&1 | head -60
+
+# TOSA 模型（同理，--iree-input-type=tosa）
+# ONNX → TOSA（iree-import-onnx 可用）
+```
+
+## 9. 关键结论
+
+1. **前端 = 框架方言 → 输入方言 → linalg**，下游（Flow→Stream→HAL→codegen→VM）完全共享。
+2. **StableHLO 和 TOSA 只是两条不同的"入口"**，最终都在 `ConvertStableHloToIreeInputDialects`
+   / `TosaToLinalgExt` 收敛到 linalg/linalg-ext。
+3. **linalg.generic 是前端到 codegen 的分水岭**：之前是"命名高层 op"，之后是"结构化计算"，
+   后续所有 tiling/vectorization 都在 linalg 上做。
@@ -0,0 +1,91 @@
+# 阶段 4a：Flow（图级优化 + dispatch 划分）✅
+
+> 2026-08-15。对应规划阶段 4 前半。源码：`src/iree/compiler/src/iree/compiler/Dialect/Flow/`。
+> IR 变换示例见 [02-architecture.md](02-architecture.md) 边界 2。
+
+## 1. Flow 的职责
+
+Flow 是 IREE 的**图级 IR**，在输入方言（linalg）和 Stream（异步）之间。两大职责：
+
+1. **图级优化**：算子融合（fusion）、reshape 折叠、类型推导、量化转换、全局常量处理等
+2. **dispatch 划分**：把图切成一个个 **dispatch region**，再 outline 成独立的 `flow.executable`
+   —— 每个 dispatch = 一个可独立编译执行的 kernel（对应 TVM 里划分出的一个 kernel）
+
+## 2. 核心 op（定义在 `IR/FlowOps.td`）
+
+| op | 作用 |
+|---|---|
+| `flow.tensor.empty/load/store/slice/reshape/update` | Flow 层的张量操作（值语义） |
+| `flow.tensor.barrier` | 张量屏障，标记 dispatch 边界（前身是 `iree_tensor_ext.compute_barrier`） |
+| `flow.tensor.transfer` | 跨设备张量传输 |
+| `flow.dispatch.region` | **dispatch region**：一个可执行单元（核心） |
+| `flow.dispatch` / `flow.dispatch.workgroups` | 调用 dispatch（启动一个 kernel） |
+| `flow.executable` / `flow.executable.export` | 可执行体 + 导出点 |
+| `flow.dispatch.workgroup_count/id/size` | workgroup 相关内建 |
+| `flow.channel.*` + `flow.collective.*` | 多设备通信（AllGather/AllReduce/AllToAll/SendRecv） |
+| `flow.parameter.load/write` | 参数（权重）处理 |
+
+## 3. 关键 pass（`Transforms/` 目录）
+
+| pass | 文件 | 作用 |
+|---|---|---|
+| `ConvertToFlowPass` | ConvertToFlow.cpp | 输入方言 → Flow |
+| `FusionPreprocessingPass` / `ElementwiseOpFusionPass` / `FuseMultiUseElementwiseProducerPass` | （dispatch creation） | **算子融合**（elementwise 融合进 reduce/matmul） |
+| `FormDispatchRegionsPass` | FormDispatchRegions.cpp | **形成 dispatch region**（核心） |
+| `CloneProducersIntoDispatchRegionsPass` | （dispatch creation） | 把 producer 克隆进 dispatch region |
+| `ConvertDispatchRegionsToWorkgroupsPass` | ConvertRegionToWorkgroups.cpp | dispatch region → workgroup 语义 |
+| `OutlineDispatchRegionsPass` | OutlineDispatchRegions.cpp | **outline 成 flow.executable** |
+| `DeduplicateExecutablesPass` | DeduplicateExecutables.cpp | 去重相同 executable |
+| `CleanupTensorShapesPass` / `CaptureDynamicDimsPass` | 同名 .cpp | 形状清理 / 动态维度捕获 |
+
+## 4. dispatch region 形成（核心概念）
+
+**dispatch region = 一个要交给 codegen 独立编译的计算单元**（对应一个 kernel 启动）。
+
+形成逻辑（`FormDispatchRegionsPass`）：
+1. 以「root op」（如 `linalg.matmul`、reduction）为锚点
+2. 把可融合的 elementwise producer 克隆进同一个 region（fusion）
+3. 用 `flow.tensor.barrier` 标记 region 的输入/输出边界
+4. 后续 outline 成 `flow.executable`
+
+IR 变化（阶段 2 已见，精简）：
+
+```mlir
+// FormDispatchRegionsPass 之后：matmul + fill 被包进 dispatch.region
+%6 = flow.dispatch.region -> (tensor<2x2xf32>) {
+  %8 = linalg.fill ins(%cst : f32) outs(%7 : tensor<2x2xf32>) -> tensor<2x2xf32>
+  %9 = linalg.matmul ins(%1, %3 : ...) outs(%8 : ...) -> tensor<2x2xf32>
+  flow.return %9 : tensor<2x2xf32>
+}
+
+// OutlineDispatchRegionsPass 之后：outline 成 executable，主函数只剩 flow.dispatch 调用
+flow.executable private @matmul_dispatch_0 { ... }   // 独立的可执行体
+%2 = flow.dispatch @matmul_dispatch_0::@matmul_dispatch_0_matmul_2x2x4_f32(%0, %1) : ...
+```
+
+## 5. 与 TVM 0.6 对照
+
+| IREE Flow | TVM 0.6 Relay |
+|---|---|
+| 图级优化（fusion/reshape/type propagation） | Relay 的 fusion / 图优化 pass |
+| dispatch region 形成 | 确定哪些算子合成一个 kernel（relay fusion 后的分组） |
+| flow.executable | 一个编译好的 kernel |
+| flow.tensor.*（值语义张量） | Relay 的 TensorType 表达式 |
+| flow.collective.* | Relay 里的通信算子（较晚版本才有） |
+
+## 6. 源码阅读路径
+
+```
+src/iree/compiler/src/iree/compiler/Dialect/Flow/
+├── IR/FlowOps.td          # op 定义（TableGen）—— 先读这个，认识所有 op
+├── IR/FlowOps.cpp         # op 实现（verifier、folder、builder）
+├── IR/FlowInterfaces.td   # 接口定义（如 Dispatchable/Partitionable）
+├── Transforms/FormDispatchRegions.cpp   # dispatch 划分核心算法
+├── Transforms/OutlineDispatchRegions.cpp
+├── Transforms/ConvertRegionToWorkgroups.cpp
+└── Conversion/TensorToFlow/  # tensor → flow 转换
+```
+
+> 学习建议：先读 `FlowOps.td` 认 op，再看 `FormDispatchRegions.cpp` 理解划分算法。
+> 划分算法在 IREE 3.x 已重构为基于「partitionable interface + workload」的方式
+> （不是老的 DispatchableOpInterface），读的时候注意版本。
@@ -0,0 +1,126 @@
+# 阶段 4b：Stream（异步调度）✅
+
+> 2026-08-15。对应规划阶段 4 后半。源码：`src/iree/compiler/src/iree/compiler/Dialect/Stream/`。
+> **这是 IREE 最有特色、也是 TVM 没有对应物的一层，值得重点理解。**
+
+## 1. Stream 的职责：把同步图变成异步执行模型
+
+Flow 层产出的还是「同步的图」（每个 op 依赖前一个的 result）。Stream 层把它变成**异步执行模型**：
+
+- 每个计算（dispatch）变成可异步提交的工作
+- 用 **timepoint**（类似 future/fence）显式表达「某批异步工作何时完成」
+- 用 **resource** 管理内存的生命周期（分配/持有/释放）
+- 用 **affinity** 标注每个 op 在哪个设备上执行
+
+这一层是 IREE 支持多设备异步调度的核心。
+
+## 2. 核心概念
+
+### (1) timepoint —— 异步同步原语（灵魂）
+
+`!stream.timepoint` = 一个「未来某个时刻会完成」的句柄，用来表达数据依赖和同步：
+
+```mlir
+// stream.async.execute 提交一批异步工作，返回 (results, timepoint)
+%results, %result_timepoint = stream.async.execute ... {
+  %4 = stream.async.dispatch @matmul_dispatch_0::@...(%arg2[...], %arg3[...]) : ...
+  stream.yield %4 : !stream.resource<...>
+} => !stream.timepoint
+
+// stream.timepoint.await 等待 timepoint 完成，取出结果
+%2 = stream.timepoint.await %result_timepoint => %results : !stream.resource<...>
+```
+
+### (2) resource —— 内存资源（带显式生命周期）
+
+tensor（值语义）在 Stream 层变成 `!stream.resource<...>`（带大小），生命周期显式管理：
+
+```mlir
+%1 = stream.tensor.import on(#hal.device.affinity<@__device_0>) %arg0 : !hal.buffer_view -> tensor<2x4xf32> in !stream.resource<external>{%0}
+```
+
+resource 有：分配（alloc）、持有（retain）、释放（release）、子视图（subview）、
+瞬态（transient）。生命周期由 **ARC（Automatic Reference Counting）** pass 自动管理。
+
+### (3) affinity —— 设备亲和性
+
+每个 op 标注 `on(#hal.device.affinity<@__device_0>)`，表示在哪个设备上执行。
+跨设备的 op 需要 `stream.async.transfer` 来传输数据。
+
+## 3. 四个抽象层次（关键！）
+
+Stream dialect 内部有从高到低的层次，编译时逐步下降：
+
+| 层次 | op 前缀 | 何时产生 | 说明 |
+|---|---|---|---|
+| tensor | `stream.tensor.*` | ConvertToStreamPass | 值语义张量（最接近 Flow） |
+| async | `stream.async.*` + `stream.timepoint.*` | ScheduleExecutionPass | 异步执行 + timepoint 同步 |
+| cmd | `stream.cmd.*` | 接近 HAL | 命令缓冲（dispatch/copy/fill/flush） |
+| resource | `stream.resource.*` | 全流程 | 内存资源生命周期 |
+
+## 4. 核心 op（`IR/StreamOps.td`）
+
+- **tensor 层**：`stream.tensor.import/export`、`stream.tensor.dispatch`、`stream.tensor.sizeof`
+- **async 层**：`stream.async.execute`、`stream.async.dispatch`、`stream.async.transfer`、
+  `stream.async.copy/fill/load/store`、`stream.async.barrier`
+- **timepoint**：`stream.timepoint.await/barrier/join/import/export/immediate`
+- **cmd 层**：`stream.cmd.dispatch/execute/copy/fill/flush/invalidate/serial/concurrent`
+- **resource**：`stream.resource.alloc/dealloca/size/retain/release/subview/transients`
+- **executable**：`stream.executable`、`stream.binding.subspan`
+
+## 5. 关键 pass（`Transforms/` + `Analysis/`）
+
+| pass | 文件 | 作用 |
+|---|---|---|
+| `ConvertToStreamPass` | ConvertToStream.cpp | Flow → Stream（核心入口） |
+| `ScheduleExecutionPass` | ScheduleExecution.cpp | **调度**：图 → async.execute + timepoint |
+| `ScheduleConcurrencyPass` | ScheduleConcurrency.cpp | 并发调度（并行 workgroup 重叠） |
+| `PropagateTimepointsPass` | PropagateTimepoints.cpp | timepoint 传播（优化同步点） |
+| `ScheduleAllocationPass` | ScheduleAllocation.cpp | 资源分配调度 |
+| `AutomaticReferenceCountingPass` | AutomaticReferenceCounting.cpp | **ARC**：自动引用计数管理资源生命周期 |
+| `ElideTimepointsPass` / `ElideAsyncCopiesPass` | 同名 | 消除冗余 timepoint / async copy |
+| `ReuseAllocationsPass` / `RefineUsagePass` | 同名 | 内存复用 / 使用细化 |
+
+**Analysis/**（调度前的分析）：
+- `Affinity.cpp` —— 设备亲和性分析
+- `Partitioning.cpp` —— 分区分析
+- `ResourceHazards.cpp` / `ResourceUsage.cpp` —— 资源冲突 / 使用分析
+
+## 6. IR 变换（阶段 2 已见，精简回顾）
+
+```mlir
+// ConvertToStreamPass：Flow → Stream，tensor 变 stream.resource，标上 affinity
+%7 = stream.tensor.dispatch on(#hal.device.affinity<@__device_0>) @matmul_dispatch_0::@...(%2, %5) : ...
+
+// ScheduleExecutionPass：包进 async.execute，产生 timepoint
+%results, %result_timepoint = stream.async.execute ... { ... } => !stream.timepoint
+%2 = stream.timepoint.await %result_timepoint => %results : !stream.resource<...>
+```
+
+## 7. 与 TVM 0.6 对照
+
+| Stream 概念 | TVM 0.6 | 说明 |
+|---|---|---|
+| timepoint / async.execute | **无直接对应** | TVM 的异步主要在 runtime（线程池/stream），不在编译期 IR 里显式建模 |
+| resource 生命周期 | **无**（TVM 的 buffer 生命周期由 runtime 管） | IREE 在编译期显式建模并 ARC |
+| affinity | device 注解 | 类似但 IREE 更显式 |
+| scheduling | 无（TVM 是同步图 + runtime 调度） | IREE 把调度变成 IR 变换 |
+
+**核心差异**：TVM 把「图是同步的、异步是 runtime 的事」；IREE 把「异步调度」提升为**编译期的显式 IR 层**，
+可以针对多设备做静态调度优化。这就是为什么 IREE 在多设备/异构场景更强。
+
+## 8. 源码阅读路径
+
+```
+src/iree/compiler/src/iree/compiler/Dialect/Stream/
+├── IR/StreamOps.td            # op 定义（先读，认识 tensor/async/cmd/resource 四层）
+├── Analysis/                  # Affinity / Partitioning / ResourceHazards / ResourceUsage
+├── Transforms/ConvertToStream.cpp      # Flow → Stream
+├── Transforms/ScheduleExecution.cpp    # 异步调度（核心）
+├── Transforms/PropagateTimepoints.cpp  # timepoint 优化
+└── Transforms/AutomaticReferenceCounting.cpp  # 资源生命周期 ARC
+```
+
+> 学习建议：先读 `ScheduleExecution.cpp` 理解「怎么从同步图变异步」，
+> 再读 `PropagateTimepoints.cpp` 理解「timepoint 怎么优化」，
+> 最后读 `AutomaticReferenceCounting.cpp` 理解「资源生命周期怎么管」。
@@ -0,0 +1,92 @@
+# 阶段 5：HAL（设备抽象）✅
+
+> 2026-08-15。对应规划阶段 5。源码：`src/iree/compiler/src/iree/compiler/Dialect/HAL/`。
+> HAL 既是「设备抽象层」，也是「codegen 的入口」（通过 TargetBackend）。
+
+## 1. HAL 的职责
+
+HAL（Hardware Abstraction Layer）把「设备」抽象成统一的概念，屏蔽 CPU/CUDA/Vulkan 等差异，
+并把 Stream 层已经调度好的 `stream.cmd.*` 命令转成 `hal.*` 设备操作。
+
+核心抽象有五个：**device、buffer/buffer_view、allocator、command_buffer、executable**（外加 channel 通信）。
+
+## 2. 核心 op（`IR/HALOps.td`）
+
+| 类别 | 关键 op | 作用 |
+|---|---|---|
+| **Device** | `hal.device.query`、`hal.device.allocator`、`hal.devices.get/count`、`hal.device.resolve` | 查询设备属性、获取设备 |
+| **Buffer** | `hal.buffer.subspan/load/store/length`、`hal.buffer.assert` | 设备内存 buffer |
+| **BufferView** | `hal.buffer_view.create`、`hal.buffer_view.buffer`、`hal.buffer_view.assert` | 带 shape/type 的 buffer 视图 |
+| **Allocator** | `hal.allocator.allocate/import/select` | 内存分配 |
+| **CommandBuffer** | `hal.command_buffer.create/dispatch/copy_buffer/fill_buffer/execution_barrier/finalize` | **记录命令**（核心） |
+| **DeviceQueue** | `hal.device.queue.execute/execute.indirect/copy/fill/flush` | 提交命令队列 |
+| **Executable** | `hal.executable`、`hal.executable.variant`、`hal.executable.export`、`hal.executable.binary` | 设备可执行代码 |
+| **Channel** | `hal.channel.create/rank_and_count/split` | 多设备通信 |
+
+## 3. Target Backend —— codegen 的入口（关键）
+
+`HAL/Target/` 定义了 **TargetBackend** 抽象（`TargetBackend.cpp/h` + `TargetRegistry.cpp`）：
+- 每个后端（LLVM-CPU、LLVM-GPU/CUDA、SPIR-V/Vulkan 等）实现一个 TargetBackend
+- 编译时按 `--iree-hal-target-backends=llvm-cpu` 选择后端
+- **可执行体的 codegen 就发生在这里**：`ConfigureExecutables` → `TranslateExecutables`
+  会调用对应 TargetBackend，把 `flow.executable` 里的 linalg 编译成设备代码
+
+> 这就是阶段 5（HAL）和阶段 6（codegen）的连接点：HAL 的 TargetBackend 是 codegen 的调度者。
+
+## 4. 关键 pass（`Transforms/` + `Conversion/`）
+
+| pass | 文件 | 作用 |
+|---|---|---|
+| `ConvertToHALPass` | ConvertToHAL.cpp | **Stream → HAL**（主入口） |
+| `AssignTargetDevicesPass` / `MaterializeTargetDevicesPass` | 同名 | 目标设备分配 |
+| `ResolveDevicePromisesPass` / `ResolveDeviceAliasesPass` | 同名 | 设备解析 |
+| `ConfigureExecutablesPass` | ConfigureExecutables.cpp | 配置可执行体（选后端、变体） |
+| `TranslateExecutablesPass` | TranslateExecutables.cpp | **调用后端做 codegen** |
+| `LinkExecutablesPass` / `SerializeExecutablesPass` | 同名 | 链接 / 序列化可执行体 |
+| `PruneExecutablesPass` / `HoistExecutableObjectsPass` | 同名 | 裁剪 / 提升可执行体对象 |
+| `ElideRedundantCommandsPass` | ElideRedundantCommands.cpp | 消除冗余命令 |
+
+**Conversion/**：
+- `StreamToHAL/` —— Stream → HAL 转换
+- `HALToVM/` —— HAL → VM 转换（阶段 7 的入口）
+- `TypeConverter.cpp` —— 类型转换
+
+## 5. IR 变换（阶段 2 已见，精简）
+
+```mlir
+// ConvertToHALPass 之后：stream.executable → hal.executable.variant
+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(...) { ... } }  // 已 codegen 成 LLVM
+  }
+}
+// #pipeline_layout 定义缓冲区绑定（2 ReadOnly + 1 输出 storage_buffer）
+```
+
+## 6. 与 TVM 0.6 对照
+
+| HAL 概念 | TVM 0.6 | 说明 |
+|---|---|---|
+| hal.device / allocator | DeviceAPI（CPU/GPU） | 设备抽象 |
+| hal.buffer / buffer_view | DLTensor / NDArray | 内存 + 元数据 |
+| hal.command_buffer | 无显式对应（TVM 直接生成 kernel 调用） | IREE 显式建模命令录制 |
+| hal.executable | 编译出的 kernel（.so / .ptx） | 设备代码 |
+| TargetBackend | codegen 的 build_module 分派 | 后端注册 |
+
+## 7. 源码阅读路径
+
+```
+src/iree/compiler/src/iree/compiler/Dialect/HAL/
+├── IR/HALOps.td                    # op 定义（先读，认识五类抽象）
+├── Target/TargetBackend.h          # 后端抽象（codegen 入口）
+├── Target/TargetRegistry.cpp       # 后端注册表
+├── Target/Devices/                 # 具体设备
+├── Transforms/ConvertToHAL.cpp     # Stream → HAL
+├── Transforms/TranslateExecutables.cpp  # 调用后端 codegen
+└── Conversion/StreamToHAL/         # 转换细节
+```
+
+> 学习建议：读 `TargetBackend.h` 理解「后端如何抽象」，这是 HAL 和 codegen 的接口。
+> `hal.imports.mlir` 列出了 HAL 模块对外暴露的所有函数（就是 VM 层 import 的那些）。
@@ -0,0 +1,92 @@
+# 阶段 6：Codegen（可执行体 lowering）✅
+
+> 2026-08-15。对应规划阶段 6。源码：`src/iree/compiler/src/iree/compiler/Codegen/`。
+> **这是你之前接触过的部分，重点是建立「MLIR pass ↔ TVM schedule」的映射。**
+
+## 1. Codegen 的职责
+
+把 HAL 层交给它的 `flow.executable`（内含 linalg 计算）lowering 成设备代码（LLVM IR / SPIR-V）。
+由 HAL 的 TargetBackend 触发（见 [07-hal.md](07-hal.md) 第 3 节）。
+
+按后端组织：
+- `LLVMCPU/` —— CPU（x86/ARM）
+- `LLVMGPU/` —— GPU（CUDA via NVVM、ROCm via ROCDL）
+- `SPIRV/` —— Vulkan
+- `VMVX/` —— 简单可移植参考后端
+- `WGSL/` —— WebGPU
+- `Common/` —— 跨后端共享逻辑
+
+## 2. 关键 lowering 链（LLVMCPU，来自阶段 2 pass 序列）
+
+```
+LLVMCPUSelectLoweringStrategy      # 选 tiling/向量化策略（对应 TVM 选 schedule）
+→ TileAndDistributeToWorkgroups   # workgroup 级 tiling
+→ LLVMCPUTileAndFuseProducerConsumer  # tile-and-fuse（对应 TVM 的 compute_at + fuse）
+→ LLVMCPUTile                     # tiling
+→ LLVMCPUPeel                     # loop peeling
+→ LLVMCPUTileToVectorSize         # tile 到向量大小
+→ GenericVectorization            # 向量化（对应 TVM 的 vectorize）
+→ IREEComprehensiveBufferize      # bufferize：tensor → memref（对应 TVM storage lowering）
+→ LLVMCPUVirtualVectorLowering    # 向量 lowering
+→ ConvertToLLVM                   # 转 LLVM IR
+```
+
+## 3. MLIR pass ↔ TVM schedule 对照（你的主场）
+
+| TVM 0.6 schedule 原语 | IREE codegen 的 MLIR pass |
+|---|---|
+| `split`（循环拆分） | `LLVMCPUTile` + `LLVMCPUTileToVectorSize` |
+| `reorder`（循环重排） | tiling pass 里隐含 |
+| `fuse` + `compute_at` | `LLVMCPUTileAndFuseProducerConsumer` |
+| `vectorize` | `GenericVectorization` |
+| `unroll` | `LLVMCPUPeel` + 向量 lowering |
+| buffer 分配 / storage | `IREEComprehensiveBufferize`（tensor→memref） |
+| codegen（LLVM/CUDA） | `ConvertToLLVM` / `ConvertToNVVM` / `ConvertToROCDL` |
+
+**核心差异**：TVM 的 schedule 是**用户手写/自动搜索**的显式变换序列；IREE 的 codegen 是
+**声明式 pass pipeline**，策略由 `LLVMCPUSelectLoweringStrategy` 基于 op 类型/shape 自动选择。
+（IREE 也有 `transform` dialect 支持手写 schedule，但那是在 linalg 之上。）
+
+## 4. Dispatch ABI（运行时↔可执行体的接口，关键）
+
+codegen 出来的 LLVM 函数（阶段 2 dump 见过）签名是：
+
+```mlir
+llvm.func @matmul_dispatch_0_matmul_2x2x4_f32(%arg0: !llvm.ptr, %arg1: !llvm.ptr, %arg2: !llvm.ptr) -> i32
+```
+
+- 参数不是直接传 buffer，而是通过 `iree_hal_executable_dispatch_state_v0_t` 结构体（`DispatchABI.cpp` 定义）
+  取 bindings —— 这是运行时和编译出的 kernel 之间的固定 ABI
+- 对应 TVM 里 kernel 的参数打包（`TVMArgs` / 显式参数）
+
+## 5. 关键文件
+
+| 文件 | 作用 |
+|---|---|
+| `LLVMCPU/LLVMCPUSelectLoweringStrategy.cpp` | 选策略（tiling 尺寸、向量化方案）——核心 |
+| `LLVMCPU/LLVMCPUTileAndFuseProducerConsumer.cpp` | tile-and-fuse |
+| `LLVMCPU/LLVMCPUTile.cpp` / `LLVMCPUPeel.cpp` | tiling / peeling |
+| `LLVMCPU/LLVMCPUTileToVectorSize.cpp` | tile 到向量 |
+| `LLVMCPU/LLVMCPUVirtualVectorLowering.cpp` | 向量 lowering |
+| `LLVMCPU/DispatchABI.cpp` | dispatch ABI |
+| `LLVMCPU/ConvertToLLVM.cpp` / `LLVMCPULowerExecutableTarget.cpp` | 转 LLVM |
+| `LLVMGPU/ConvertToNVVM.cpp` / `ConvertToROCDL.cpp` | GPU 后端 |
+| `LLVMGPU/KernelConfig.cpp` | GPU kernel 配置 |
+
+## 6. 源码阅读路径
+
+```
+src/iree/compiler/src/iree/compiler/Codegen/
+├── Common/               # 跨后端共享
+├── Interfaces/           # 接口
+├── LLVMCPU/              # CPU 后端
+│   ├── LLVMCPUSelectLoweringStrategy.cpp   # 先读这个（策略选择）
+│   ├── LLVMCPUTileAndFuseProducerConsumer.cpp
+│   └── DispatchABI.cpp
+├── LLVMGPU/              # GPU 后端（NVVM/ROCDL）
+├── SPIRV/                # Vulkan
+└── Transforms/           # 通用变换
+```
+
+> 学习建议：你有 codegen 经验，重点看 `LLVMCPUSelectLoweringStrategy.cpp` 理解「策略怎么选」，
+> 以及 `DispatchABI.cpp` 理解「kernel 和运行时怎么对接」。GPU 后端看 `LLVMGPU/KernelConfig.cpp`。
@@ -0,0 +1,123 @@
+# 阶段 7：运行时（VM + iree-runtime）✅
+
+> 2026-08-15。对应规划阶段 7。分两部分：
+> **编译侧** `src/iree/compiler/src/iree/compiler/Dialect/VM/`（生成字节码）
+> **运行时** `src/iree/runtime/src/iree/`（C 实现）。
+
+## 1. 整体图景
+
+编译的最后一步是把 HAL 模块序列化成 **VM 字节码**（`.vmfb`），运行时加载并解释执行：
+
+```
+HAL 模块 → VM 转换（HALToVM）→ vm.module（字节码 IR）→ 序列化 → .vmfb 文件
+                                                              ↓
+                                          iree-runtime 加载 → VM 解释器 → 调 hal.* import → HAL driver 执行
+```
+
+## 2. VM = 紧凑字节码虚拟机（编译侧）
+
+VM dialect 定义了一套**字节码级的 IR**，核心 op（`IR/VMOps.td`）：
+
+| op | 作用 |
+|---|---|
+| `vm.module` | 顶层字节码模块 |
+| `vm.func` | VM 函数 |
+| `vm.import` / `vm.export` | 导入外部函数（如 `hal.command_buffer.dispatch`）/ 导出 |
+| `vm.call` / `vm.call.variadic` / `vm.call.yieldable` | 调用（含可让出调用，协作式调度） |
+| `vm.branch` / `vm.cond_branch` / `vm.branch_table` | 字节码级控制流 |
+| `vm.global.load/store.ref` | 全局引用变量 |
+| `vm.const.ref.rodata/zero` | 常量 |
+| `vm.list.*` / `vm.buffer.*` | 列表 / buffer 操作 |
+| `vm.discard.refs` | **引用计数回收**（显式 GC） |
+
+### 字节码格式
+
+用 **FlatBuffers** 定义（`runtime/src/iree/schemas/bytecode_module_def.fbs`），
+是紧凑、可直接 mmap 加载的二进制格式。阶段 2 用 `iree-dump-module` 看到的就是反汇编：
+
+```
+Required Types: i32, i64, !hal.allocator, !hal.buffer, !hal.buffer_view,
+                !hal.command_buffer, !hal.device, !hal.executable, !hal.fence, !vm.buffer
+Module Dependencies: hal, version >= 6
+Imported Functions:
+  [ 7] hal.command_buffer.dispatch(...) -> ()
+  [11] hal.device.queue.execute.indirect(...) -> ()
+  [16] hal.fence.await(i32, i64, tuple<!vm.ref<?>>...) -> (i32)
+```
+
+## 3. iree-runtime（C 运行时）
+
+`runtime/src/iree/` 目录：
+
+| 目录 | 作用 |
+|---|---|
+| `vm/` + `vm/bytecode/` | **VM 解释器**（字节码解码 + 分发执行） |
+| `hal/` + `hal/drivers/` | HAL 运行时 + 各设备驱动 |
+| `task/` | 多线程任务调度系统（local-task driver 的基础） |
+| `base/` | 基础工具（status、allocator、logging） |
+| `builtins/` / `modules/` | VM 内置函数 / 内置模块（如 hal 模块） |
+| `schemas/` | FlatBuffers schema（字节码 + 各后端可执行体格式） |
+| `runtime/` | 顶层 API（`iree_runtime_instance` 等） |
+| `io/` | 参数/文件 IO |
+
+### HAL drivers（真正的设备后端）
+
+`hal/drivers/` 下每个目录是一个设备实现：
+
+| driver | 设备 |
+|---|---|
+| `local_task` | **CPU 多线程**（基于 task 系统）—— 默认 |
+| `local_sync` | CPU 同步（调试用） |
+| `cuda` / `hip` / `amdgpu` | NVIDIA / AMD GPU |
+| `vulkan` | Vulkan |
+| `metal` | Apple Metal |
+| `webgpu` | WebGPU |
+| `null` | 空设备（测试） |
+
+> 当前环境只装了 base 包，运行时只有 local-task/local-sync。CUDA driver 需要额外构建。
+
+## 4. 与 TVM 0.6 对照
+
+| IREE | TVM 0.6 |
+|---|---|
+| VM 字节码（.vmfb） | graph runtime / relay VM 的序列化 |
+| vm.import（hal.*） | 生成的 kernel 调用 |
+| HAL driver（local-task/cuda/vulkan） | DeviceAPI（CPU/CUDA/OpenCL/Metal） |
+| task 系统 | CPU 线程池 runtime |
+| 引用计数 GC | 无（TVM 用 DLTensor 引用计数） |
+
+**核心差异**：IREE 的 VM 是**完整字节码虚拟机**（带控制流、全局变量、yieldable 调用），
+而 TVM 的 graph runtime 只是「按拓扑顺序执行节点」。IREE 的 VM 更接近一个真正的 VM。
+
+## 5. 源码阅读路径
+
+```
+# 编译侧（字节码生成）
+src/iree/compiler/src/iree/compiler/Dialect/VM/
+├── IR/VMOps.td              # VM op 定义
+├── Conversion/HALToVM/      # HAL → VM 转换
+└── Target/                  # 序列化到字节码
+
+# 运行时（C）
+src/iree/runtime/src/iree/
+├── vm/                      # VM 解释器
+├── hal/drivers/local_task/  # CPU 多线程 driver（默认）
+├── hal/drivers/cuda/        # CUDA driver
+├── task/                    # 任务调度
+└── schemas/bytecode_module_def.fbs  # 字节码格式
+```
+
+## 6. 全流程回顾（7 个阶段串起来）
+
+```
+输入(StableHLO/TOSA) → [前端] → linalg → [Flow] dispatch 划分 → [Stream] 异步调度
+→ [HAL] 设备抽象 + 触发 codegen → [Codegen] linalg→LLVM/SPIR-V → [VM] 字节码 → [runtime] 执行
+```
+
+你现在已经把 IREE 从输入到运行时**完整走了一遍**。每个阶段都有对应笔记和源码路径。
+
+## 下一步（阶段 8 实践，可选）
+
+- 写一个自定义 pass（如自定义 fusion 规则），需要源码构建
+- 配 CUDA 后端，看 GPU 的 codegen（LLVMGPU → NVVM）
+- 深入某一层读源码（建议 Stream 或 Codegen）
@@ -0,0 +1,119 @@
+# 从源码构建 IREE（写 pass 的前置）
+
+> 2026-08-15。为了写自定义 pass，需要从源码构建 IREE 编译器（pip 包只有二进制，没有头文件和构建系统）。
+> 记录完整步骤，方便复现。
+
+## 1. 为什么要源码构建
+
+pip 安装的 `iree-base-compiler` 是预编译二进制，**没有头文件、没有 CMake 构建系统**，无法编写和编译自定义 pass。
+写 pass 需要：源码 + LLVM/MLIR 头文件 + 构建系统。
+
+## 2. 前提（工具链）
+
+```bash
+sudo apt install clang-15 lld-15      # 编译器 + 链接器（Ubuntu 22.04 最新可用 clang）
+.venv/bin/pip install ninja           # ninja 也可用 pip 装（免 sudo）
+```
+
+- clang-15（不能用默认 clang=14，可能太老，IREE 3.11 用 LLVM 23）
+- lld-15（快速链接，可选但推荐）
+- ninja（构建生成器）
+- cmake（系统已有）
+
+## 3. 子模块
+
+```bash
+# 关键子模块：llvm-project（IREE fork，~1.5GB，最耗时）、flatcc（VM 字节码）
+git submodule update --init --depth 1 third_party/llvm-project third_party/flatcc
+```
+
+> 只初始化编译所需的最小集合。写 pass 测试用 linalg 输入，不需要 stablehlo/torch-mlir 等子模块。
+
+## 4. 配置（cmake）
+
+```bash
+cmake -G Ninja -B build/ -S . \
+  -DCMAKE_BUILD_TYPE=RelWithDebInfo \
+  -DCMAKE_C_COMPILER=clang-15 -DCMAKE_CXX_COMPILER=clang++-15 \
+  -DIREE_ENABLE_LLD=ON \
+  -DIREE_BUILD_TESTS=OFF -DIREE_BUILD_SAMPLES=OFF -DIREE_BUILD_PYTHON_BINDINGS=OFF \
+  -DIREE_TARGET_BACKEND_DEFAULTS=OFF -DIREE_TARGET_BACKEND_LLVM_CPU=ON \
+  -DIREE_HAL_DRIVER_DEFAULTS=OFF -DIREE_HAL_DRIVER_LOCAL_SYNC=ON -DIREE_HAL_DRIVER_LOCAL_TASK=ON
+```
+
+关键选项：
+- `-DIREE_BUILD_TESTS/SAMPLES/PYTHON_BINDINGS=OFF` —— 大幅缩减构建（只建编译器）
+- `-DIREE_TARGET_BACKEND_DEFAULTS=OFF + LLVM_CPU=ON` —— 只保留 CPU 后端
+- `-DIREE_HAL_DRIVER_DEFAULTS=OFF + LOCAL_SYNC/TASK=ON` —— 只保留 CPU 运行时驱动
+- `-DCMAKE_BUILD_TYPE=RelWithDebInfo` —— 有调试信息 + 优化（写 pass 推荐）
+
+## 5. 构建
+
+```bash
+# 只构建编译器工具（iree-compile + iree-opt），比全量快很多
+cmake --build build/ --target iree-compile iree-opt -j 16
+```
+
+- 结果在 `build/tools/iree-compile`、`build/tools/iree-opt`
+- `iree-opt` 是测试 pass 的关键工具（单 pass 运行，无需完整编译管线）
+- 首次全量构建 ~30-60 分钟（16 核）
+
+## 6. 验证
+
+```bash
+build/tools/iree-compile --version
+build/tools/iree-opt --help
+```
+
+## 7. 写 pass 的流程（后续）
+
+1. 在某个 dialect 的 `Transforms/` 下写 pass 的 `.cpp`
+2. 在对应 `Passes.td` 里注册 pass（TableGen）
+3. 在 `CMakeLists.txt` 里加入源文件
+4. 重新构建（增量构建，快）
+5. 用 `iree-opt --pass-name` 单测 pass
+
+## 8. 备选：gcc 构建（不用 sudo）
+
+如果没有 sudo 装 clang，可以用系统 gcc 11.4（支持 C++17）：
+```bash
+cmake -G Ninja -B build/ -S . -DCMAKE_BUILD_TYPE=RelWithDebInfo \
+  -DIREE_BUILD_TESTS=OFF -DIREE_BUILD_SAMPLES=OFF -DIREE_BUILD_PYTHON_BINDINGS=OFF \
+  -DIREE_TARGET_BACKEND_DEFAULTS=OFF -DIREE_TARGET_BACKEND_LLVM_CPU=ON \
+  -DIREE_HAL_DRIVER_DEFAULTS=OFF -DIREE_HAL_DRIVER_LOCAL_SYNC=ON -DIREE_HAL_DRIVER_LOCAL_TASK=ON
+```
+（gcc 是 IREE 官方"best effort"支持，能用但稍慢。）
+
+## 9. 实际踩坑记录（本次构建的关键修复）
+
+### 9.1 工具链版本
+- **cmake 3.22 太旧**（需 3.26+）→ `pip install cmake`（4.4.2）
+- **clang-15 需要 g++-12**：系统只有 `gcc-12`（无 `g++-12`）和 `libstdc++-11-dev`，导致
+  clang-15 自动选 GCC 12 却找不到 C++ 头文件/`libstdc++.so`。
+  修复：`sudo apt install g++-12`（带 `libstdc++-12-dev`）
+- **lld**：装了 `lld-15` 但没有无版本号 `ld.lld`。修复：软链接
+  `ln -sf /usr/bin/ld.lld-15 .venv/bin/ld.lld`，并把 `.venv/bin` 加入 PATH
+
+### 9.2 llvm-project 子模块用 tarball 而非 git
+git 方式下载 llvm-project（1.5GB）在 VPN 下被 GitHub 限速（~30KB/s），改用 **codeload tarball**（280MB 压缩）+ 断点续传，快很多。但 git 不认 tarball 解压出的目录为子模块，导致 IREE 的子模块检查失败。
+
+修复：把 `build_tools/scripts/git/check_submodule_init.py` 改成 no-op（本地构建，手动提供源码）。
+
+### 9.3 只保留必要组件（大幅缩减构建）
+
+最终配置命令：
+```bash
+cmake -G Ninja -B build/ -S . \
+  -DCMAKE_BUILD_TYPE=RelWithDebInfo \
+  -DCMAKE_C_COMPILER=clang-15 -DCMAKE_CXX_COMPILER=clang++-15 \
+  -DIREE_ENABLE_LLD=ON \
+  -DIREE_BUILD_TESTS=OFF -DIREE_BUILD_SAMPLES=OFF -DIREE_BUILD_PYTHON_BINDINGS=OFF -DIREE_BUILD_BENCHMARKS=OFF \
+  -DIREE_TARGET_BACKEND_DEFAULTS=OFF -DIREE_TARGET_BACKEND_LLVM_CPU=ON \
+  -DIREE_HAL_DRIVER_DEFAULTS=OFF -DIREE_HAL_DRIVER_LOCAL_SYNC=ON -DIREE_HAL_DRIVER_LOCAL_TASK=ON \
+  -DIREE_INPUT_STABLEHLO=OFF -DIREE_INPUT_TORCH=OFF -DIREE_INPUT_TOSA=OFF
+```
+关键：`-DIREE_BUILD_BENCHMARKS=OFF`（否则要 benchmark 子模块）、
+`-DIREE_INPUT_*=OFF`（否则要 stablehlo/torch-mlir 子模块）。
+
+### 9.4 写 pass 最小依赖子模块
+只用 **llvm-project + flatcc** 两个子模块（linalg 输入 + VM 字节码输出），其余全关。
@@ -0,0 +1,181 @@
+# 在 IREE 里写一个自定义 Pass
+
+> 2026-08-16。基于源码构建的 iree-compile，写 pass 的完整指南。
+> 以在 **Flow dialect** 里加一个最小 pass 为例，代码可直接复制。
+
+## 1. Pass 的三件套
+
+IREE 的 pass 需要改 3 个地方（都在某个 dialect 的 `Transforms/` 下）：
+
+| 文件 | 作用 |
+|---|---|
+| `Passes.td` | TableGen 声明 pass（名字、作用的操作类型、选项） |
+| `HelloWorld.cpp`（新） | pass 实现（`runOnOperation`） |
+| `CMakeLists.txt` | 把新 .cpp 加入编译 |
+
+以 Flow 为例，路径是：
+`src/iree/compiler/src/iree/compiler/Dialect/Flow/Transforms/`
+
+## 2. 最小示例：HelloWorld pass（打印所有 op）
+
+### 2.1 在 `Passes.td` 里声明（加一行）
+
+在 `Passes.td` 末尾加：
+```td
+def HelloWorldPass :
+    Pass<"iree-flow-hello-world", "mlir::ModuleOp"> {
+  let summary = "Prints all operations in the module (learning example).";
+}
+```
+
+- `iree-flow-hello-world` = pass 的 CLI 名字（`iree-opt --iree-flow-hello-world` 用）
+- `mlir::ModuleOp` = 这个 pass 作用于整个 module（`getOperation()` 返回 `ModuleOp`）
+- 传 `""` 则作用于任意 op（泛型 pass）
+
+### 2.2 新建 `HelloWorld.cpp`
+
+```cpp
+#include "iree/compiler/Dialect/Flow/Transforms/Passes.h"
+
+#include "mlir/IR/BuiltinOps.h"
+#include "mlir/Pass/Pass.h"
+#include "llvm/Support/raw_ostream.h"
+
+namespace mlir::iree_compiler::IREE::Flow {
+
+#define GEN_PASS_DEF_HELLOWORLDPASS
+#include "iree/compiler/Dialect/Flow/Transforms/Passes.h.inc"
+
+namespace {
+
+struct HelloWorldPass : impl::HelloWorldPassBase<HelloWorldPass> {
+  using impl::HelloWorldPassBase<HelloWorldPass>::HelloWorldPassBase;
+
+  void runOnOperation() override {
+    // 遍历 module 里所有 op 并打印名字
+    getOperation().walk([&](Operation *op) {
+      llvm::errs() << "  [op] " << op->getName() << "\n";
+    });
+  }
+};
+
+} // namespace
+} // namespace mlir::iree_compiler::IREE::Flow
+```
+
+### 2.3 在 `CMakeLists.txt` 里加入
+
+在 `Transforms/CMakeLists.txt` 的 `SRCS` 列表（按字母序）加：
+```cmake
+"HelloWorld.cpp"
+```
+
+### 2.4 重新构建
+
+```bash
+export PATH="/home/ylhan/ai/ai_compiler/.venv/bin:$PATH"
+cmake --build /home/ylhan/ai/ai_compiler/src/iree-build/ --target iree-opt -j 16
+```
+（只重建 `iree-opt`，增量编译很快，几秒到几十秒）
+
+## 3. 关键结构解析（理解这套模板）
+
+```cpp
+#define GEN_PASS_DEF_HELLOWORLDPASS
+#include ".../Passes.h.inc"
+```
+- TableGen 从 `Passes.td` 生成 `Passes.h.inc`，里面定义了 `impl::HelloWorldPassBase<HelloWorldPass>` 基类
+- `GEN_PASS_DEF_*` 宏告诉 `.inc` 文件「我要生成这个 pass 的基类定义」
+
+```cpp
+struct HelloWorldPass : impl::HelloWorldPassBase<HelloWorldPass> { ... }
+```
+- 你的 pass 继承生成的基类（CRTP 模式），`impl::` 是 `IREE::Flow::impl` 命名空间
+
+**不需要手动注册** —— `Passes.cpp` 里的 `registerPasses()` 是从 TableGen 生成的，
+会自动 `PassRegistration` 所有在 `Passes.td` 里声明的 pass。
+
+## 4. 用 `iree-opt` 单测 pass
+
+`iree-opt` 是测试 pass 的核心工具（不跑完整编译管线，只跑你指定的 pass）：
+
+```bash
+# 在构建目录里
+build/tools/iree-opt --iree-flow-hello-world input.mlir
+```
+
+先写一个测试输入 `test.mlir`：
+```mlir
+func.func @main(%arg0: tensor<2x4xf32>, %arg1: tensor<4x2xf32>) -> tensor<2x2xf32> {
+  %c0 = arith.constant 0.0 : f32
+  %empty = tensor.empty() : tensor<2x2xf32>
+  %init = linalg.fill ins(%c0 : f32) outs(%empty : tensor<2x2xf32>) -> tensor<2x2xf32>
+  %r = linalg.matmul ins(%arg0, %arg1 : tensor<2x4xf32>, tensor<4x2xf32>)
+                     outs(%init : tensor<2x2xf32>) -> tensor<2x2xf32>
+  return %r : tensor<2x2xf32>
+}
+```
+
+预期输出：打印 module 里所有 op 的名字（`func.func`, `arith.constant`, `tensor.empty`, `linalg.fill`, `linalg.matmul`, ...）。
+
+### 常用调试 flag
+
+```bash
+iree-opt --iree-flow-hello-world --mlir-print-ir-after-all test.mlir   # 看 pass 后的 IR
+iree-opt --pass-pipeline='builtin.module(iree-flow-hello-world)' test.mlir  # 显式指定 pipeline
+```
+
+## 5. 进阶：真正做变换的 pass（Pattern Rewriting）
+
+HelloWorld 只是「看」，真正的 pass 是「改」。核心是 **Pattern Rewriting**（见 [03-mlir-basics.md](03-mlir-basics.md) 第 7 节）：
+
+```cpp
+// 一个 pattern：把 arith.addi %x, 0 简化成 %x
+struct SimplifyAddZero : OpRewritePattern<arith::AddIOp> {
+  using OpRewritePattern<arith::AddIOp>::OpRewritePattern;
+
+  LogicalResult matchAndRewrite(arith::AddIOp op,
+                                PatternRewriter &rewriter) const override {
+    if (auto cst = op.getRhs().getDefiningOp<arith::ConstantOp>()) {
+      if (cst.value() == 0) {
+        rewriter.replaceOp(op, op.getLhs());
+        return success();
+      }
+    }
+    return failure();
+  }
+};
+
+struct MySimplifyPass : impl::MySimplifyPassBase<MySimplifyPass> {
+  void runOnOperation() override {
+    RewritePatternSet patterns(&getContext());
+    patterns.add<SimplifyAddZero>(&getContext());
+    if (failed(applyPatternsGreedily(getOperation(), std::move(patterns)))) {
+      return signalPassFailure();
+    }
+  }
+};
+```
+
+- `OpRewritePattern<OpTy>`：声明「匹配某个 op，重写」
+- `matchAndRewrite`：匹配逻辑 + 重写逻辑，`rewriter.replaceOp` 替换 op
+- `applyPatternsGreedily`：贪婪遍历 IR 应用 patterns（对应 IREE 里的 `CanonicalizePass`）
+
+## 6. 下一步方向
+
+1. 先把 HelloWorld 跑通，理解「声明 → 实现 → 构建 → 单测」的闭环
+2. 改写一个真正的 pattern（如上）
+3. 深入读 IREE 现有的 pass 源码，照着改：
+   - 简单 pattern：`Flow/Transforms/Canonicalize.cpp`
+   - 复杂 pass：`Codegen/LLVMCPU/LLVMCPUTile.cpp`
+4. 写 pass 时配合 `--mlir-print-ir-before-all` / `--mlir-print-ir-after-all` 看 IR 变化
+
+## 7. 参考：pass 位置速查
+
+| 想改的层 | 路径 |
+|---|---|
+| Flow（图级） | `compiler/src/iree/compiler/Dialect/Flow/Transforms/` |
+| Stream（异步） | `compiler/src/iree/compiler/Dialect/Stream/Transforms/` |
+| HAL（设备） | `compiler/src/iree/compiler/Dialect/HAL/Transforms/` |
+| Codegen（CPU） | `compiler/src/iree/compiler/Codegen/LLVMCPU/` |
+| Codegen（GPU） | `compiler/src/iree/compiler/Codegen/LLVMGPU/` |
@@ -0,0 +1,38 @@
+# IREE 学习文档
+
+> IREE 系统学习笔记。目标：在 TVM 0.6 经验基础上，**系统掌握 IREE 整体**（前端输入 →
+> 编译中段 → HAL → codegen → 运行时），而非仅 codegen。
+
+## 目录结构
+
+- [00-plan.md](00-plan.md) —— 学习规划（含「是否先学 MLIR」的结论、阶段划分、资源）
+- 后续各阶段笔记按 `NN-主题.md` 命名，边学边补充到下方进度表
+
+## 学习进度
+
+| 阶段 | 主题 | 状态 |
+|---|---|---|
+| 0 | 环境准备（编译 / 安装 IREE） | ✅ 完成（[01-env.md](01-env.md)） |
+| 1 | MLIR 基础（IREE 全流程的地基） | ✅ 速查已整理（[03-mlir-basics.md](03-mlir-basics.md)） |
+| 2 | IREE 架构总览与完整管线 | ✅ 完成（[02-architecture.md](02-architecture.md)） |
+| 3 | 前端输入（StableHLO / TOSA） | ✅ 完成（[04-frontend.md](04-frontend.md)） |
+| 4 | 编译中段（Flow + Stream） | ✅ 完成（[05-flow.md](05-flow.md)、[06-stream.md](06-stream.md)） |
+| 5 | HAL 设备抽象 | ✅ 完成（[07-hal.md](07-hal.md)） |
+| 6 | Codegen（可执行体 lowering） | ✅ 完成（[08-codegen.md](08-codegen.md)） |
+| 7 | 运行时（VM + iree-runtime） | ✅ 完成（[09-runtime.md](09-runtime.md)） |
+| 8 | 实践项目 | ✅ 完成（源码构建 + 自定义 pass，[10-build-from-source.md](10-build-from-source.md)、[11-writing-a-pass.md](11-writing-a-pass.md)） |
+
+## 个人背景
+
+- 熟悉 TVM 0.6：Relay、TE/TIR、schedule 原语、codegen（LLVM/CUDA/OpenCL）、runtime
+- 接触过 IREE 的 codegen 部分
+- 需求：系统补齐 IREE 整体，并建立 TVM → IREE/MLIR 的映射
+
+## 关键映射（TVM 0.6 → IREE/MLIR）
+
+| TVM 0.6 | IREE / MLIR |
+|---|---|
+| Relay（图级 IR、fusion） | Flow dialect + StableHLO/TOSA |
+| TE / TIR + schedule 原语 | Linalg + Vector + tiling/vectorization |
+| codegen（LLVM/CUDA/OpenCL） | 翻译到 LLVM / SPIR-V |
+| runtime（graph runtime / VM） | HAL + VM |