一、为什么编译器需要多级中间表示
传统编译器(LLVM IR)只有一个中间表示层级,难以同时表达高层语义(如张量并行策略、分布式通信)和底层硬件细节(如 SIMD 指令、共享内存别名)。当 AI 框架遇到多样化的硬件后端(CPU/GPU/TPU/NPU)时,一种 IR 适配所有目标的模式导致了大量不可复用、不可组合的 ad-hoc lowering pass。
MLIR(Multi-Level Intermediate Representation)是 Google/LLVM 社区为解决这一挑战提出的编译器基础设施,其核心创新不在于发明新 IR,而在于提供一套可组合的 IR 抽象机制,让不同层级的表示可以在同一个编译流水线中共存和转换。
二、MLIR 核心概念:Dialect 与 Operation
MLIR 用 Dialect 定义一组语义相关的 Operation、Type 和 Attribute。每个 Dialect 抽象一个计算层级的语义:
- std / arith:标量算术运算
- scf / cf:结构化控制流(for/if)、分支控制流
- tensor / linalg:张量和结构化线性代数运算
- memref:内存引用操作
- gpu / nvvm / rocdl:GPU 后端特定指令
- func:函数抽象
- vector:SIMD 向量运算
- affine:多面体表示,用于循环变换
2.1 Operation:MLIR 的最小语义单元
// Tensor 级别的 matmul 在 MLIR 中的表示
%c = linalg.matmul
ins(%a, %b : tensor<128x256xf32>, tensor<256x512xf32>)
outs(%init : tensor<128x512xf32>)
-> tensor<128x512xf32>
关键:linalg.matmul 在 tensor Dialect 中是一个纯语义操作,它没有被绑定到任何具体硬件实现。编译器的整个工作就是将这个高层语义逐步 lowering 到硬件可执行的指令序列。
三、渐进式 Lowering:从框架语义到机器码
MLIR 的核心设计哲学是渐进式 lowering:每次只转换一层语义,保留其他层级的优化机会。以一个 LLM 推理中 FFN 层的前向计算为例:
// Level 0: 高层框架语义 (torch-mlir / tf-mlir)
torch.aten.mm %input, %weight -> %1_out
torch.aten.gelu %1_out -> %2_out
// Level 1: Tensor 语义 (linalg Dialect)
linalg.matmul ins(%input, %weight) outs(%init) -> tensor
linalg.elemwise_unary "gelu"(%gemm_out) -> tensor
// Level 2: 循环 + 缓冲化 (loop & bufferization)
scf.for %i = 0 to B step 1 {
scf.for %j = 0 to S step 1 {
scf.for %k = 0 to H step 1 {
%a = memref.load [%input, %i, %k]
%b = memref.load [%weight, %k, %j]
%sum = arith.addf %partial, %product
memref.store %sum, [%output, %i, %j]
}
}
}
// Level 3: GPU 并行化 (gpu Dialect)
gpu.launch blocks(%bx, %by) threads(%tx, %ty) {
%tid = gpu.thread_id
%gid = gpu.block_id
gpu.barrier
}
// Level 4: PTX 指令 (NVVM Dialect)
nvvm.mma.sync.aligned.m8n8k16.row.col.f16.f16.f16.f32
四、MLIR 的关键优化 Pass 体系
MLIR 的优化过程由一系列 pass 组成,每个 pass 在特定 Dialect 层级进行变换。以下是高性能推理编译器中的典型通路:
4.1 高层优化(Linalg 层级)
- Fusion Pass:将 element-wise 操作融合进 matmul 的循环体,减少内存带宽压力
- Tiling Pass:将大矩阵分块,适配缓存层次和共享内存大小
- Vectorization Pass:将内层循环展开为 SIMD vector 操作
- In-place Bufferization:将 tensor 值语义转换为 memref 引用语义,消除不必要的内存分配
4.2 GPU 特定优化(GPU / NVVM 层级)
- Copy To Shared Memory:将 global memory 数据预取到 on-chip shared memory
- Warp-level Matrix Multiply (WMMAt):映射到 Tensor Core / Matrix Engine 硬件指令
- Kernel Launch Fusion:合并多个小 kernel 以减少启动开销
- Register Tiling:平衡寄存器压力与指令级并行度
4.3 Flash Attention 的 MLIR 实现
Flash Attention 算法的 tiling 策略利用 online softmax 将 O(N^2) 内存复杂度降到 O(N),在 MLIR 中可以非常自然地表达:
scf.for %kv_block = 0 to N_KV step B_kv {
%k_tile = memref.subview %K[%kv_block, 0] [B_kv, D]
%v_tile = memref.subview %V[%kv_block, 0] [B_kv, D]
scf.for %q_block = 0 to N_Q step B_q {
%partial = @matmul(%Q_block, %k_tile)
%row_max = @row_max_reduced(%partial)
%corrected = @rescale_by_exp(%partial, %row_max)
%prob = @softmax(%corrected)
%O_block = @update_output(%O_block, %prob, %v_tile)
}
}
五、IREE 与 MLIR 在生产推理部署中的应用
IREE(Intermediate Representation Execution Environment)是 Google 基于 MLIR 构建的端到端 AI 部署编译器栈:
- 输入:TensorFlow/PyTorch 模型通过 torch-mlir / tf-mlir 转化为高层 Dialect
- 编译流:高层 Dialect 经层级 lowering 到 HAL(Hardware Abstraction Layer)再到 Vulkan/Metal/LLVM
- 输出:可直接在目标设备上运行的静态/动态库
IREE 的核心优势在于:将编译时优化(如算子融合、布局变换、tuning search)提前到 build phase,大幅减少运行时开销。在手机上运行 Stable Diffusion 级别的推理任务时,IREE 通常比 eager execution 快 3-5 倍。
5.1 与 TVM/XLA 的对比
TVM 使用 Relay IR + TE schedule 抽象,XLA 使用 HLO IR + GpuCompiler pass。它们的问题在于中间表示难以扩展和组合。MLIR 通过 Dialect 和 Op 接口实现了更好的组件化:TVM 的 schedule 是一种全局变换策略(与具体 Op 耦合),而 MLIR 的 pass 可以在同一 IR 中灵活组合,使跨 Dialect 的优化(如 tiling 同时应用于 linalg 和 tensor)成为可能。
六、MLIR 在 RISC-V 扩展指令集开发中的应用
MLIR 正在扩展到通用计算领域。RISC-V 向量扩展的后端正在采用 MLIR lowering 流程。利用 MLIR 的向量 Dialect(vector Dialect),RISC-V 向量指令集的自动代码生成与向量长度无关编程模型(VLA, Vector Length Agnostic)达到了较好的一致性,为 RISC-V 高性能生态提供了编译器基础设施支持。
七、小结
MLIR 的价值不在于某个特定优化 pass 的先进性,而在于其可组合、可扩展的基础设施设计。Dialect 抽象让不同层级(张量/循环/GPU/目标指令)的语义保持独立,而 pass 管理器可以灵活组合这些层级之间的 lowering 与优化。
对于构建自有 AI 编译器栈的团队(如各大厂的推理引擎团队),MLIR 提供了一个比 TVM 更灵活、比 XLA 更易扩展的中间层。未来两年,我们预计看到更多基于 MLIR 的端到端异构编译解决方案从实验室走向生产环境。

发表评论 取消回复