15 AI 编译器与 IR 栈洞察

章节编号:15
所属层:C 编译与算子层(核心壁垒;承接 07/08 软硬接口)
关联 ADR:ADR-025(编译框架)、ADR-026(IR 分层与 DSA Dialect)、ADR-027(代码生成 vs 库)、ADR-028(前端与可扩展性)、ADR-033(LLM-Native 编译器开发)
上游依赖:01(具身模型/算子)、02(数据流/精度)、03(内存/KV)、07(ISA/编译兼容/MLIR 托管)、08(同步/显式 DMA)

学习目标

前置知识:读过 07/08 章(知道自研 DSA「软件管理内存 + 显式 DMA + 静态调度」的语义约束),对 PyTorch/ONNX 模型的计算图有概念,听过「编译器把高层算子降低成机器指令」这句话即可。无需 MLIR 或 LLVM 源码经验。
学完产出:① 能画出「框架模型 → 高层 IR(TOSA/StableHLO)→ Linalg → 自定义 DSA Dialect → 后端」的多层 lowering 链条,并说清每一层负责哪一类变换、为什么这种分层能分离关注点;② 能解释 Triton→Linalg 迁移层如何让「一份 Triton kernel 复用到自研 DSA」,把算子移植成本从 O(算子数 × 硬件数)降到近似 O(算子数);③ 能说清 IREE 的 structured codegen 为什么适合落地 DSA——AOT 编译为 .vmfb、HAL 抽象设备、绕过运行时直连加速器这三件事各解决什么;④ 能对比「代码生成为主」与「算子库依赖」两条路线,理解 XLA/TVM「抽象过度」的教训;⑤ 能读懂本章 ADR-025~028 的取舍逻辑,判断一个新 VLA 结构该走「自定义算子扩展」还是「现有 Dialect lowering」。
阅读姿势:盯住一条主线——「编译器的每一层 IR,都是在一个恰当的抽象高度上,把一类优化做干净后再往下交棒」。多层 lowering 不是叠床架屋,而是让 tiling/fusion 在 Linalg 层做、内存与 DMA 编排在 DSA Dialect 层做、指令选择在后端做,各层互不干扰——这就是「分离关注点」在编译栈里的具体形态,也是 DSA + 全栈 MLIR 路线成败的关键。

1. 范围与目标

编译器是把 PyTorch/ONNX 模型变成芯片高效可执行的"中枢",也是 DSA + 全栈 MLIR + 编译兼容路线下性能与生态的成败点。本章定前端导入、IR 分层、中端优化、编译框架与 kernel 前端(含 Triton)选型;后端 Codegen 见 16 章。

核心问题

自研编译器还是复用 MLIR/IREE?
IR 如何分层?自定义 DSA Dialect 放在哪一层?
以代码生成为主还是依赖算子库?是否引入 Triton 等可移植 kernel DSL 作为前端?
前端如何对接 PyTorch/ONNX,并保留对 VLA 演进的可扩展性?

2. 需求洞察(承接上游)

编译兼容(07 章) ⇒ 每代硬件靠重编译适配 → 编译器必须强、可自动生成高性能代码。
软件管理内存 + 显式 DMA + 静态调度(08 章) ⇒ 编译器要能表达并自动编排 scratchpad/DMA/同步(难点)。
batch≈1(01 章) ⇒ 强算子融合 + tiling降访存、提小批量利用率。
VLA 快速演进(01 章) ⇒ IR/编译器须可扩展,避免 XLA 式"抽象过度"导致新结构跑不动。
混合精度/块缩放(02 章)+ 量化(19 章) ⇒ 量化感知 lowering。
可驾驭 + AI-Native ⇒ 最大化复用开源 + LLM 辅助 pass 开发。

2b. 多层 IR 与逐层 lowering:分离关注点的编译栈骨架

在展开各家方案前,先把「多层 IR」这个底座讲透——它是理解后文所有取舍的地基。一个 AI 编译器不是「模型进、指令出」的黑盒,而是一串抽象高度递减的 IR(中间表示),每往下降一层就把一类优化做干净。这套逐层 lowering 的价值,一句话概括是分离关注点(separation of concerns):图级融合、循环 tiling、内存/DMA 编排、指令选择,各在自己最合适的抽象层做,互不污染。

逐层读这张图——每一层「只操心一件事」:

层	抽象高度	本层负责的优化	为什么放这一层
前端导入	最高(框架算子)	把 `aten.*` / ONNX op 映射成标准 dialect	与框架版本解耦,一次接入长期复用
高层 IR(StableHLO/TOSA)	整图语义	算子等价改写、常量折叠、shape 推断	此层信息最全,做图级重写代价最低
Linalg(结构化中层)	循环 + 张量	tiling / fusion / layout	Linalg 基于 memref strided view,地址表达式可预测,且跨 tiling/fusion 保留高层结构信息——避免 BLAS pattern 被过早破坏
DSA Dialect(自研)	硬件数据流	scratchpad 分配、显式 DMA、set/wait 同步	08 章的软件管理内存语义只有到这一层才「有对象可编排」
后端(16 章)	指令	指令选择、寄存器分配、调度	复用 LLVM 成熟基础设施

为什么分层能分离关注点:若把 tiling 和 DMA 编排混在同一层做,一个 pass 就要同时理解「循环结构」和「片上缓冲地址」,复杂度爆炸且难以复用。分层后,Linalg 层的 tiling pass 完全不需要知道下游是 GPU 还是 DSA;DSA Dialect 层的 DMA 编排也不关心上游模型是 VLA 还是 CNN——每一层的 pass 只依赖本层的语义契约,可独立开发、独立测试、独立替换。这正是 MLIR「可复用方言 + 渐进式 lowering(progressive lowering)」设计哲学的落点,也是本章坚持复用 MLIR/IREE 而非从零造的根本理由(相关思考见「深入思考·思考题 1」)。

3. 技术现状与趋势(点名 + 来源)

3.1 MLIR:多层 IR 基础设施(我们的底座)

Linalg dialect:数据空间 tiling(基于 memref strided view,地址表达式可预测)+ 跨 tiling/fusion 保留高层信息(避免 BLAS pattern 被破坏的相位排序难题);配合 Tensor/Affine/Vector/Bufferization(one-shot)。
前端:TOSA / StableHLO / torch-mlir / onnx-mlir 把框架模型导入 MLIR。
为什么 Linalg 是「腰」:它高到还能表达「这是一个 matmul / conv」的结构(不像降到 for 循环后 pattern 尽失),又低到能承接 tiling/fusion/layout 这些循环级变换——框架前端、Triton 前端、乃至手写算子,最终都汇聚到 Linalg 这个统一入口,再往下分流到各家 DSA Dialect。这个「沙漏腰」的位置,是后文 Triton→Linalg 迁移层与 DSA lowering 都锚定在 Linalg 的原因。
判断:MLIR 是事实标准的多层 IR 基座,直接复用(契合 07 章 MLIR 托管)。

3.2 IREE:端到端 MLIR 编译器 + 运行时

MLIR-based 端到端编译 + 运行时,从数据中心到边缘;嵌入式二进制可低至 30KB;AOT 编译为 .vmfb 制品;HAL 统一设备抽象(内存/dispatch/同步);后端 llvm-cpu/vulkan-spirv/cuda/rocm/vmvx;嵌入式可绕过运行时直接对接自定义加速器;IREE 是 structured codegen 的主要使用者。
案例:Roofline 为 IREE 建 Vulkan 1.3 HAL,实现边缘异构异步执行 + >4GB 权重。
structured codegen 是什么:IREE 不把张量算子直接砸成标量循环,而是沿 Linalg 的结构化算子(structured op)逐步 tiling/distribute/vectorize,全程保留「这是矩阵块运算」的结构语义,直到最后一步才降到目标后端。这套「结构化」路径对 DSA 尤其关键——DSA 的价值恰恰在于用矩阵/向量专用单元吃下结构化块运算,而 structured codegen 天然把块结构一路带到 HAL dispatch 边界(如何落地见「深入思考·思考题 3」)。
判断:复用 IREE 作为编译+运行时框架(关联 12 章 Runtime),新增自定义 DSA HAL 后端 + DSA Dialect,而非从零造。

3.3 代码生成 vs 算子库;XLA/TVM 的教训

库依赖派(XLA/TorchInductor/TensorRT):重度依赖 cuDNN/cuBLAS 做计算密集算子;XLA/TVM 第二代"抽象过度",GenAI 时代需自定义 kernel/动态行为时被绕过(Pallas/Triton)。
全代码生成派(PolyBlocks、TPP/libxsmm):编译器生成全部底层代码,单算子可与 cuDNN/cuBLAS/手写 kernel 媲美;MLIR 高层 tiling/fusion 更强的跨算子优化。
判断:走代码生成为主 + 手写算子兜底(02 章/组织"算子归编译")+ 可选 Triton 前端补生态(取舍见 3.5.1);避免 XLA 抽象过度——保留自定义算子/扩展路径适配 VLA。

3.4 趋势

业界收敛到MLIR 编排高层变换(tiling/fusion/layout)+ 底层代码生成/微内核的混合;不绑死单一手调库。

3.5 主流 AI 推理芯片的编译方案对比(优劣,2025–2026 最新)

2025–2026 选型判断:Triton 成为跨厂商可移植 kernel DSL,各家普遍"编程层 Ascend C/CUDA 等自有 + Triton 入口;IR 层 MLIR;后端 LLVM 或自研"。

公司/芯片	编程层(最新)	IR / 编译路径	后端	优势 / 劣势
NVIDIA	CUDA + Triton(+ CUTLASS/cuDNN/TensorRT)	Triton IR → PTX 或 CUDA Tile IR(MLIR,Blackwell)	NVPTX(LLVM)/ tileiras	生态碾压、Tile IR 新引入 / 库依赖、闭源、绑 CUDA
Google TPU	JAX + Pallas/Mosaic;XLA/OpenXLA	StableHLO → XLA;Pallas→Mosaic(MLIR)	XLA 自有 codegen	图融合强、编译兼容 / XLA 抽象过度→Pallas 绕过
AMD	HIP + Triton(+ Composable Kernel)	Triton/MLIR	AMDGPU(LLVM)	开放、CUDA 可移植 / 生态追赶
华为昇腾	Ascend C(C++)+ Triton(triton-ascend)	Triton IR → Linalg → AscendNPU IR(MLIR)	毕昇 BiSheng(LLVM 基)→ CCE intrinsics	软硬协同、开放 AscendNPU IR + 拥抱 Triton/FlagTree、国产可控 / 生态 building(注:早期 TBE 为历史方案,AKG 非主流)
寒武纪	BANG C + Triton	MLIR / Triton	CNCC(LLVM 基)	自有栈 + Triton 接入 / 生态小
地平线	天工开物(自动编译器主导)	自有编译	自有	利用率/能效高、车规 / 可编程性低(新算子上不了 BPU)
Groq	GroqFlow(MLIR)	MLIR + 周期级静态调度	自研后端	确定性、零抖动 / 仅 curated、封闭
摩尔线程/沐曦/天数等	MUSA/CUDA 兼容 + Triton(FlagTree)	MLIR	LLVM 基	CUDA/Triton 双兼容降迁移 / 能效追赶
Tesla	全私有	私有	私有	极致协同 / 封闭无生态

横向判断:① IR 层全面收敛 MLIR(NVIDIA Tile IR、Pallas Mosaic、AscendNPU IR、Groq GroqFlow、寒武纪/摩尔线程皆 MLIR);② 后端多为 LLVM 或其 fork(NVPTX、AMDGPU、华为毕昇、寒武纪 CNCC),纯自研后端集中在数据流/确定性芯片(TPU、Groq);③ 编程层正以 Triton 为可移植 DSL 统一入口(见 3.5.1);④ 可编程性仍是分水岭(地平线封闭 vs NVIDIA 开放)。

3.5.1 Triton 作为 kernel 前端:采纳分析(中立洞察)

事实(2025–2026):Triton 已成跨厂商可移植 kernel DSL——NVIDIA(PTX + CUDA Tile IR 新后端,MLIR 基)、AMD、Intel XPU、华为昇腾(triton-ascend:Triton IR→Linalg→AscendNPU IR)、寒武纪、摩尔线程、沐曦、天数、海光等;FlagTree(单仓多后端 Triton,FlagOpen 开源体系)+ FlagGems(Triton 算子库)推动统一。
迁移层的关键机制(Triton→Linalg):一份 Triton kernel(@triton.jit 修饰的 Python)先被编译成 Triton IR,triton-ascend 这类路径不直接下到某家 GPU 后端,而是先 lower 到 Linalg——落在前文 3.1 那个「沙漏腰」上。落到 Linalg 之后,后续 tiling/fusion/DSA lowering 全部复用我们已有的 MLIR 通路。这就把「N 个算子 × M 种硬件」的移植矩阵,拆成「N 个算子 → Linalg」+「Linalg → M 种硬件」两段:算子作者只写一次 Triton,硬件方只维护一条 Linalg→DSA 通路(量化收益见「深入思考·思考题 2」)。
采纳的收益:triton-ascend 路径 Triton IR → Linalg → 厂商 MLIR Dialect 与我们 IR 分层(Linalg → DSA Dialect)同构 → 接入边际成本低(复用 Triton→Linalg),可直接复用 vLLM/SGLang/FlagGems 等现成 Triton 算子生态,补生态短板,并自然承接 07 章"CUDA/Triton 迁移层"。
不采纳/弱采纳的理由:我们自研 MLIR codegen + 手写算子已覆盖核心性能路径,Triton 并非性能必需;引入 Triton 需持续跟随其上游演进、维护前端适配,且面向我们自研 DSA 的 Triton 算子调优仍需投入;对确定性硬实时控制路径,Triton 的 GPU 式编程模型并非最佳。
中立结论:Triton 定位为"可选的可移植 kernel 前端 / 生态迁移层"(锦上添花,非核心依赖)——核心性能与确定性仍由自研 MLIR codegen + 手写算子保证;是否投入、投入多深,取决于客户迁移需求与生态 ROI(见待决项),按实际需求决策,不预设必须或必不使用。

3.6 LLM-Native 软件栈与 LLM/Agent 辅助编译器开发(最新趋势 + 路径选型)

这是与"业界成熟方案"并列的第二层思考:用 LLM/Agent 重构编译器/算子的开发方式(AI for Compilers),契合我们 AI-Native 组织。

业界最新趋势(2025–2026,点名 + 来源)

KernelBench(Stanford,250 题)成为"LLM 能否写高效 GPU kernel"的标准基准;早期 LLM 仅匹配约 20% 编译器水平。
Agentic kernel 生成快速突破:
KernelAgent(PyTorch):多智能体 + 硬件信号引导闭环,250 题 100% 正确,H100 上达 89% roofline,较 torch.compile 1.56×。
KernelFalcon(PyTorch):deep agents 分层分解,首个 250 题 L1/L2/L3 全 100% 正确;并行探索 + 执行验证。
CUDA Agent(agentic RL):KernelBench L1/L2/L3 比 torch.compile 快 100%/100%/92%,超 Claude Opus 4.5 / Gemini 3 Pro。
AutoKernel / GEAK(AMD MI300):Amdahl 排序瓶颈 + 六层优化 playbook,Triton/CUDA 双后端。
共性范式:Profile(Nsight 等)→ 诊断瓶颈 → 生成 kernel/pass → 执行验证(数值+性能)→ 迭代;基于真实硬件信号(非静态启发式),在"长尾"(异常算子/动态 shape/异构融合)上超过传统编译器。
关键判断:LLM/Agent 不取代编译器,而是增强——补长尾、生成候选、降人力;正确性靠执行验证闭环、性能靠真实 profiling 反馈。

结合 LLM/Agent 的编译器高效开发路径(我们的选型)

LLM 生成 MLIR pass / DSA Dialect lowering 样板 + 单测:把模式化的 pass/转换/测试交给 LLM,资深工程师审校——直接降编译器开发人力(关联组织 AI-Native)。
Agent 生成算子/微内核候选 + auto-tuning 搜索筛选(对标 KernelAgent/AutoKernel,关联 16/17/18 章):Profile→诊断→生成→执行验证→迭代闭环,grounded 在我们的 ISS/FPGA/硅后真实信号(而非 GPU)。
RAG + Agent CI 闭环:ISA 手册/架构 spec/寄存器表做 RAG,Agent 编排"生成→编译→跑测试→回填修复"(关联 30 章 AI-Native 平台)。
(远期)RL 训练专用 kernel-Agent 模型(CUDA Agent 式),针对我们 DSA 的 kernel 数据合成 + 自动验证奖励。

技术选型要点:私有代码大模型 + 内部 KernelBench 式基准 + 执行验证(ISS/FPGA) + profiling 反馈 + MLIR pass 生成 + Agent 编排框架;面向我们自研 DSA(非 GPU),验证闭环是成败关键。风险:LLM 生成需强验证(正确性)+ 真实硬件信号(早期 ISS 精度有限);低资源语言(自研 ISA)数据稀缺,需自建数据合成(关联 30 章)。

3b. 由演进反推的诉求(定性)

维度	当前	未来(VLA/多模型演进)	编译器含义
算子覆盖	核心算子集	新结构(MoE/世界模型)	可扩展 IR + 自定义算子路径
性能	自动生成 + 手写兜底	自动生成覆盖 80%+	强 tiling/fusion + auto-tuning(16 章)
兼容	每代重编译	频繁	编译器稳健 + CI(30 章)
内存	scratchpad/DMA 编排	KV-cache/长上下文	显式内存编排 + 量化感知

具身时效补注(2026-07):π0.7、GR00T N1.7 已进入公开可用(GA)阶段,流匹配 + RTC 异步 chunk 成为动作头事实范式(详见 M 层 21 章)。对编译器的直接含义:迭代去噪循环(while/多步)必须在 IR 中表达为 first-class loop,单步蒸馏(OneDP/OFP)则要求同一算子既能走多步 loop、也能编译成 1-NFE 快路径——这条「多步 loop + 单步快路径」双通路诉求,是 Linalg→DSA lowering 必须预留的扩展点。相关模型/频率靶标以官方发布为准。

4. 候选方案与对比矩阵

4.1 编译框架

候选	性能上限	复用/生态	风险	可驾驭度	演进性	小结
A. 从零自研编译器	4	1	5	1	3	人力/风险不可承受
B. MLIR + IREE 复用 + 自定义 DSA Dialect/后端	5	5	3	4	5	建议方案:复用前端/HAL/codegen,聚焦差异化
C. 复用 TVM	3	3	3	3	2	GenAI 适应性弱、抽象过度

4.2 代码生成策略

代码生成为主 + 手写算子兜底(高层 tiling/fusion/layout 自动,底层 codegen 见 16 章);Triton 作为可选前端(3.5.1)。

4.3 IR 分层

TOSA/StableHLO/Linalg(复用)→ Tensor/Affine/Vector/Bufferization(复用)→ 自定义 DSA Dialect(自研壁垒) → 后端(16 章)。

建议方向:B + 代码生成为主 + 上述 IR 分层。

5. 关键权衡、风险与依赖

复用 vs 控制力:复用 IREE/MLIR 降人力但需深度参与上游、跟随演进;自定义 DSA Dialect/后端是不可外包壁垒。
自动生成性能达标:性能主要压在自研 codegen + auto-tuning(16 章);若不达标需手写算子兜底(预算与 02/组织一致)——最大风险。
抽象过度陷阱:学 XLA 抽象过度会让 VLA 新结构跑不动 → 必须保留自定义算子/扩展路径。
MLIR/IREE 人才稀缺:与组织 A 组 lead 招聘强绑定。
依赖:07/08(ISA/内存语义)、16(后端)、17(算子)、18(auto-tuning)、19(量化)、12(IREE Runtime)、30(CI/AI-Native)。

6. 结论与待决项

初步结论

复用 MLIR + IREE,新增自定义 DSA Dialect + HAL 后端;不从零造、不用 TVM。核心性能由自研 MLIR codegen + 手写算子兜底保证;Triton 作为可选可移植 kernel 前端/迁移层(Triton IR→Linalg→DSA Dialect,对标 triton-ascend),按生态 ROI 决策(详见 3.5.1)。
IR 分层:TOSA/StableHLO/Linalg → 通用中端 → DSA Dialect → 后端;代码生成为主 + 手写算子兜底;Linalg 作为框架/Triton 的统一入口。
强 tiling/fusion/layout/bufferization + 量化感知 lowering;显式 scratchpad/DMA 编排(16 章实现)。
保留自定义算子/扩展路径,规避 XLA 抽象过度,适配 VLA 演进。

待决项

DSA Dialect 抽象边界(暴露多少硬件细节给上层)——与 02/07 联合定。
上游 IREE 跟随 vs fork 策略(版本管理/回合上游)——与 30 章 CI 联合定。
Triton 前端采纳程度(承接 07 待决项 + 3.5.1):仅作迁移层 vs 作为一等 kernel 前端;Triton→Linalg→DSA Dialect 的工程投入与 FlagGems/vLLM 算子复用收益评估。

ADR 候选

ADR-025 编译框架:MLIR + IREE 复用 + 自定义 DSA Dialect/后端。
ADR-026 IR 分层:TOSA/StableHLO/Linalg → 通用中端 → DSA Dialect。
ADR-027 代码生成策略:核心走自研 MLIR 代码生成 + 手写算子兜底;Triton 作为可选 kernel 前端/迁移层(Triton IR→Linalg→DSA Dialect),按生态 ROI 与客户迁移需求决定投入程度。
ADR-028 前端/可扩展性:对接 PyTorch/ONNX,保留自定义算子路径,避免抽象过度。
ADR-033 LLM-Native 编译器开发:用 LLM 生成 MLIR pass/lowering 样板 + Agent 生成算子候选(Profile→生成→执行验证→迭代,grounded 在 ISS/FPGA)+ RAG/Agent CI;建内部 KernelBench 式基准与执行验证闭环;LLM 增强而非取代编译器。

7. 端→边→中心演进影响

同一 MLIR/IREE + DSA Dialect 栈可对端/边/中心不同后端实现复用(编译兼容,07 章);差异在 codegen 配置与 HAL 后端。
不可逆点:DSA Dialect/后端接口需与 ISA 协同稳定,预留扩展。
roadmap 提示:上游 MLIR/IREE/Triton 均处高速演进,本章版本判断(triton-ascend 路径、CUDA Tile IR、FlagTree 多后端)以官方为准,选型细节随上游 release 校准。

深入思考

下面三题每题先给题干,再用 <details> 折叠一份图文并茂的参考答案。建议先合上答案自己想 3 分钟,再展开对照。

思考题 1:多层 IR 逐层 lowering 为何能分离关注点

有人质疑「一层 IR 直接从 PyTorch 降到硬件指令不是更快吗?为什么要 StableHLO→Linalg→DSA Dialect→后端这么多层」。结合 2b 节的分层骨架,论证:多层 lowering 相比「单层大编译器」在优化正确性、pass 复用、可扩展性三个维度各带来什么;并说明「tiling 放 Linalg、DMA 编排放 DSA Dialect」这条切分线为什么不能对调。

展开参考答案(含分层职责隔离图 + 单层 vs 多层对比表)

结论：每一层 IR 都在一个恰当的抽象高度上把一类优化做干净,pass 只依赖本层的语义契约,于是图级重写、循环 tiling、DMA 编排、指令选择互不污染;单层大编译器则要在同一份代码里同时理解所有抽象,复杂度与耦合度都会失控,新硬件、新算子每来一个都要动全身。

用一次「加一款新硬件」的工作量算一遍(数量级估算):

假设整栈有 约 40 个优化 pass,其中图级重写 ~10 个、tiling/fusion ~15 个、内存/DMA ~10 个、指令选择 ~5 个。
多层架构下,新增一款 DSA 只需重写「DMA 编排 + 指令选择」两层,即 ~10 + ~5 = 约 15 个 pass;上游 ~25 个图级/循环级 pass 原样复用(它们不知道下游是谁)。
单层大编译器下,图重写与 tiling 逻辑与硬件细节纠缠在同一份代码里,换硬件几乎要审视全部 ~40 个 pass——工作量约 2.7 倍,且回归风险扩散到本可不动的图级优化。
正确性维度:多层里,Linalg 的 tiling pass 有明确输入输出契约(结构化算子进、结构化算子出),可单独写 lit 测试;单层里 tiling 的正确性与 DMA 状态耦合,难以隔离验证。

为什么切分线不能对调:tiling 是循环/迭代空间的变换,它需要的信息是「这个 matmul 的 M/N/K 维度和数据布局」——这些在 Linalg 层最完整;而 DMA 编排需要的是「片上 buffer 有多大、从哪个地址搬到哪个地址、何时 set/wait」——这些语义在 Linalg 层根本不存在(Linalg 还没决定内存分配),必须等 lowering 到 DSA Dialect、bufferization 之后才「有对象可编排」。若强行把 DMA 编排提到 Linalg 层,就得在还没内存布局时凭空假设地址,假设一旦被后续 tiling 推翻就全错;反之把 tiling 压到 DSA 层,则结构化信息早已丢失,只剩裸循环,matmul pattern 认不出来,优化无从下手。分层的本质,就是让每个 pass 在「它所需信息恰好齐备」的那一层运行。

维度	单层大编译器	多层逐层 lowering
优化正确性	优化与硬件状态耦合,难隔离验证	每层有契约,可独立写 lit/单测
pass 复用	换硬件几乎重写全部	上游图级/循环级 pass 原样复用
可扩展性	新算子/新硬件牵一发动全身	新硬件只加下游层,新算子只加 Dialect op
切分依据	无明确边界	「信息恰好齐备」的抽象高度

回到正文:这正是 2b 节强调的「分离关注点」,也是本章坚持复用 MLIR/IREE、把差异化聚焦在 DSA Dialect 一层的根本理由(见 3.1「沙漏腰」与 4.1 候选 B)。

思考题 2:Triton→Linalg 迁移层如何降低算子移植成本

客户手上有一批为 NVIDIA 写的 Triton 算子(vLLM/SGLang/FlagGems 生态),现在要迁到我们自研 DSA。结合 3.5.1 节的 triton-ascend 路径,论证:为什么「Triton IR → Linalg → DSA Dialect」这条迁移层能把移植成本从「乘法级」降到「加法级」;并说明这条路径的边际成本主要落在哪里、什么情况下它反而不划算。

展开参考答案(含移植矩阵拆解图 + 成本算一遍)

结论：一份 Triton kernel 编译成 Triton IR 后,先降到 Linalg 这个厂商无关的「沙漏腰」,再复用统一的 Linalg→DSA 通路;于是「N 个算子 × M 种硬件」的乘法级移植矩阵,被拆成「N 个算子编译到 Linalg」加「一条 Linalg→DSA 通路」的加法级工作——算子作者只写一次 Triton,硬件方只维护一条 lowering,谁都不必为对方的规模买单。

用具体数字算一遍(数量级估算,假设要迁 50 个 Triton 算子、支持 3 款自研 DSA 变体):

无迁移层(手写):每个算子在每款硬件上都要手写并调优一版 kernel → 50 × 3 = 150 份手写 kernel。这是「乘法级」——算子数和硬件数相乘。
有迁移层(Triton→Linalg→DSA):
- Triton→Linalg 由上游 triton-ascend 式前端负责,50 个算子编译一次即得 Linalg(边际成本 ≈ 跑通编译 + 修少量不支持的 pattern);
- Linalg→DSA 通路是每款硬件一条,3 款 = 3 条通路,且这 3 条通路本就是自研 codegen 的主线、并非为 Triton 额外新建。
- 于是增量工作 ≈ 「50 个算子过一遍编译 + 补边角」+ 「3 条已有通路」= 加法级,而非 150 份手写。
量级对比:假设手写 + 调优一个 DSA kernel 平均 3 人日,150 份 = 450 人日;迁移层路径下,50 个算子过编译并修 pattern 平均 0.5 人日、外加 3 条通路的一次性适配各 ~20 人日,合计 ≈ 50 × 0.5 + 60 = 约 85 人日——约 5 倍 的人力压缩,且未来再加算子边际成本趋近于零。

边际成本主要落在哪:① Triton→Linalg 前端对某些 pattern(不规则 mask、动态 shape、指针算术)支持不全,需补 lowering 或回退手写;② 迁移过来的 kernel 正确但未必最优——Triton 的 tile 划分是照着 GPU SM 调的,落到 DSA 的 Cube/向量单元后,tiling/layout 常需再调优(这部分工作无法靠迁移层免除,见 16/18 章 auto-tuning)。

什么情况下反而不划算:若某算子是核心热点(如主 GEMM/Attention),它的性能直接决定端到端延迟,那么无论走不走 Triton,最终都要手写 + 深度调优——此时迁移层只省了「从零起步」,省不了「逼近峰值」那段最贵的工作。所以本章把 Triton 定位成「生态迁移层 / 长尾补位」(3.5.1 中立结论),核心热点仍由自研 MLIR codegen + 手写算子兜底——迁移层降的是长尾算子的覆盖成本,不是热点算子的调优成本。

思考题 3:IREE structured codegen 如何落地自研 DSA

你决定复用 IREE 而非从零造运行时。结合 3.2 节的 AOT .vmfb、HAL、structured codegen、绕过运行时直连加速器,论证:这四件事分别帮你把「一个编译好的 VLA 模型」落到自研 DSA 的哪一步;并说明为什么 structured codegen(保留结构化算子)比「早早降成标量循环」更适合 DSA 的矩阵/向量专用单元。

展开参考答案(含 IREE 落地 DSA 数据流图 + structured vs 标量对比表)

结论：AOT 把模型固化成 .vmfb 制品(编译期定死、运行期零 JIT,契合 08 章静态调度);HAL 用统一设备接口抽象掉「内存/dispatch/同步」的差异,让你只需实现一个自定义 DSA HAL 后端就能接管调度;structured codegen 全程保留「这是矩阵块运算」的结构语义,直到 HAL dispatch 边界才落地,正好把块结构喂给 DSA 的 Cube/向量单元;而嵌入式绕过运行时,则让极端边缘场景下 DSA 直接吃编译产物、省掉运行时开销。

四件事各落到哪一步:

AOT .vmfb:把整个模型提前编译成一个自包含制品,运行期不做即时编译、不做 kernel autotune——这正对上 08 章「静态调度」的硬约束,也让端侧无需带 JIT 编译器(嵌入式可低至 30KB)。
HAL(硬件抽象层):IREE 把「设备内存分配、命令 dispatch、跨命令同步」抽象成统一接口。落地 DSA = 实现一个自定义 DSA HAL 后端,把这些接口映射到 DSA 的 DMA 与 set/wait 原语——你只写这一层,上游编译流水与运行时调度全部复用。
structured codegen:见下面对比。
绕过运行时:对确定性硬实时或极小 footprint 场景,IREE 允许 DSA 直接对接编译产物、跳过通用运行时,进一步压延迟与内存(关联 12 章 Runtime 边界)。

为什么 structured codegen 比「早降标量循环」更适合 DSA:

维度	早早降成标量循环	structured codegen(保留结构化算子)
块结构信息	tiling 后 `matmul` 语义丢失,只剩三重 for 循环	全程保留「M×N×K 块运算」结构,到 HAL 边界才落地
专用单元映射	编译器要从裸循环「反推」出这是矩阵乘,才能映射 Cube 单元——极易失败	结构化算子可直接映射 DSA 的 16×16 矩阵指令 / 向量指令
DMA/tiling 协同	循环边界与访存耦合,难在正确抽象层编排 DMA	tiling 与 distribute 在结构化层完成,DMA 编排信息齐备
量化/layout	低精度块缩放信息在标量循环里难表达	结构化算子可携带 layout / 量化元信息下传(02/19 章)

DSA 的全部价值,在于用矩阵/向量专用单元一拍吃下一整块结构化运算(对标昇腾 Cube 一拍算 16×16×16、GPU Tensor Core 一条 MMA 吃一个块)。若编译器早早把张量算子拆成标量三重循环,专用单元就「无块可吃」,只能退化成逐元素算,DSA 的红利荡然无存。structured codegen 的意义,就是把「块」这个 DSA 最需要的结构,一路护送到硬件指令边界——这也是 3.2 节把 IREE 称作「structured codegen 的主要使用者」并选它作框架的核心原因。

回到正文:这条落地链印证了 4.1 候选 B(复用 IREE + 自定义 DSA HAL/Dialect)与 6 章结论 1——差异化只压在「DSA Dialect + HAL 后端」一层,其余全栈复用。

附:信息来源

区分公开/估计;参考时间 2026-06。

MLIR / Linalg:MLIR 官方《Linalg Dialect Rationale》(数据空间 tiling、跨变换保留高层信息);MLIR Tensor/Affine/Vector/Bufferization 文档。[公开]
IREE:iree.dev 与 iree-org/iree(MLIR-based 端到端编译+运行时、30KB 嵌入式、.vmfb、HAL、structured codegen、绕过运行时对接自定义加速器);Roofline Vulkan 1.3 HAL 案例(边缘异步异构、>4GB 权重)。[公开]
代码生成 vs 库 / XLA·TVM 教训:Modular《Democratizing AI Compute Part 6》(XLA/TVM 抽象过度、GenAI 被绕过);PolyBlocks(arXiv 编号 2603.06731 年代可疑,待核;全代码生成媲美 cuDNN/cuBLAS);TPP/libxsmm(arXiv:2404.15204,Linalg→XSMM 微内核)。[公开]
主流推理芯片编译方案(3.5,2025–2026 最新):华为昇腾——昇腾社区《CANN 全新升级:毕昇编译器开放 AscendNPU IR,全面支持 Triton》(2025-05-29)、Ascend/AscendNPU-IR 与 triton-ascend 仓(Triton IR→Linalg→AscendNPU IR(MLIR)→毕昇 BiSheng→CCE intrinsics;CANN 8.5/triton-ascend 3.2.0,2026-01);Triton 多后端——NVIDIA《CUDA Tile IR Backend for OpenAI Triton》+ Triton-to-tile-IR 仓(CUDA 13.1/Blackwell)、Intel intel-xpu-backend-for-triton、FlagTree(FlagOpen 开源体系,单仓多后端:NVIDIA/AMD/昇腾/寒武纪/摩尔线程/沐曦/天数/海光);Google Pallas/Mosaic;AMD HIP/Triton。[公开]
LLM-Native / LLM-Agent 编译(3.6):KernelBench(arXiv:2502.10517,250 题、早期约 20%);KernelAgent(PyTorch 博客,250 题 100%、H100 89% roofline、1.56× torch.compile);KernelFalcon(PyTorch 博客,deep agents 全 100%);CUDA Agent(2026 arXiv 编号待核,agentic RL,超 torch.compile 100/100/92%、超 Claude Opus 4.5/Gemini 3);AutoKernel(2026 arXiv 编号待核);GEAK(AMD MI300)。[公开]
框架/分层/代码生成/LLM-Native 取舍:基于上述公开资料的工程判断。[估计]

下一篇 → 16 后端代码生成与算子实现：把视野从「IR 分层与前端」下沉到「Linalg/DSA Dialect 如何真正落成高性能机器指令」，讲清 tiling/fusion 的具体 codegen、auto-tuning 与手写算子兜底的边界。

学习目标​

1. 范围与目标​

2. 需求洞察(承接上游)​

2b. 多层 IR 与逐层 lowering:分离关注点的编译栈骨架​

3. 技术现状与趋势(点名 + 来源)​

3.1 MLIR:多层 IR 基础设施(我们的底座)​

3.2 IREE:端到端 MLIR 编译器 + 运行时​

3.3 代码生成 vs 算子库;XLA/TVM 的教训​

3.4 趋势​

3.5 主流 AI 推理芯片的编译方案对比(优劣,2025–2026 最新)​

3.5.1 Triton 作为 kernel 前端:采纳分析(中立洞察)​

3.6 LLM-Native 软件栈与 LLM/Agent 辅助编译器开发(最新趋势 + 路径选型)​

3b. 由演进反推的诉求(定性)​

4. 候选方案与对比矩阵​

4.1 编译框架​

4.2 代码生成策略​

4.3 IR 分层​

5. 关键权衡、风险与依赖​

6. 结论与待决项​

初步结论​

待决项​

ADR 候选​

7. 端→边→中心演进影响​

深入思考​

思考题 1:多层 IR 逐层 lowering 为何能分离关注点​

思考题 2:Triton→Linalg 迁移层如何降低算子移植成本​

思考题 3:IREE structured codegen 如何落地自研 DSA​

延伸阅读​

附:信息来源​