AI 协作说明

本文由作者主导选题、技术判断与终稿审核；AI 工具协助完成资料检索、信息归纳、结构梳理与文字润色。文中观点、事实核验与引用准确性由作者负责。

资料核对时间：2026 年 9 月 23 日。 本文以官方文档、项目仓库、原始论文和课程主页为主要依据。工具链能力是这一时间点的资料快照，实际使用时应固定版本；学习优先级与选型建议是本文的判断。容量、延迟和分块数字均为分析示例，不代表设备实测。

一个大模型在 PyTorch 中生成了正确的结果，距离在端侧神经网络处理器（Neural Processing Unit，NPU）上高效执行，还有一系列问题需要解决：怎样表达随解码更新的状态，怎样把高层算子转换成硬件能接受的计算，怎样安排有限的片上存储，以及怎样让编译结果与设备运行时正确配合。

这些问题共同构成大模型编译工程。它既包含 IR 和编译变换，也包含量化表示、内存布局、设备接口与验证方法。对技术架构师而言，关键是识别每一层能决定什么、需要向下一层传递什么，以及一个局部优化能否改善完整生成过程。

本文与《大模型推理与部署工程图谱：技术栈、开源项目与学习路线》组成同一系列。上一篇建立推理与部署全景；本篇沿着模型语义 → 图与 IR → 编译优化 → 硬件映射 → 有状态执行 → 验证交付展开，把端侧 NPU 的约束放进每一步。

阅读建议：第一至四节建立全景与发展脉络，第五至七节理解编译核心，第八至十节比较工具链与运行边界，第十一、十二节学习选型和验证，第十三节执行 8 周课程与实践路线，第十四节按问题深读资料。想先决定学习投入，可先看第一、八、十三节。窄屏上的多列表格可以横向滚动。

一、先建立全景：大模型编译究竟要掌握什么

1.1 从数学运算走到可执行程序

机器学习编译（machine learning compilation） 将模型程序转换为适合目标环境执行的形式，同时保持约定的计算语义。其输出可能是机器码、设备命令、厂商图对象、序列化上下文，或包含多个子图和执行计划的模型包。

因此，“模型转换成功”“编译成功”“设备执行成功”“端到端更快”是四个需要分别验证的结果。编译也可能发生在不同时间：提前编译（Ahead-of-Time Compilation，AOT） 在运行前完成目标编译，常用于发布或部署准备阶段；即时编译（Just-in-Time Compilation，JIT） 在运行期间按需生成可执行代码或设备可加载的编译产物。编译之后仍需执行程序，才能得到推理结果。端侧可以采用两者组合。LiteRT 的 NPU 集成说明就同时提供 AOT 和设备端编译路径。

部署关注完整系统怎样交付，推理关注输入怎样经过模型生成结果，编译关注计算描述怎样转换为执行程序。三者在量化、形状、内存和状态上紧密衔接，工具链也经常同时覆盖这些环节。本篇以程序转换和硬件映射为观察重点：既检查产物本身，也回到实际推理验证它的效果。

本文主要讨论自回归大语言模型（Large Language Model，LLM）及其视觉语言模型（Vision-Language Model，VLM）扩展，设备包括移动片上系统（System-on-Chip，SoC）、嵌入式计算平台和带 NPU 的个人电脑。微控制器上的 TinyML、数据中心加速器和 FPGA 只在帮助比较机制时出现，具体容量与软件接口需要分别判断。

1.2 三种学习深度

编译方向应把 IR、程序变换、形状与内存提升为主要学习对象。 推理服务中的路由、集群调度和多机通信，在这条路线中通常只需先理解接口；它们是否成为重点，取决于实际系统边界。

1.3 三条工作主线，共享同一套基础

熟练使用厂商转换工具，可以完成很有价值的模型接入工作。若目标进一步深入编译器，就需要选择具有源码、IR 输出和测试入口的实现，练习修改与验证。两种能力可以在同一个项目中逐步衔接。

二、端侧 NPU 的约束，怎样改变编译问题

2.1 NPU 是一类硬件，具体能力来自芯片与软件组合

不同 NPU 的矩阵单元、向量单元、片上 SRAM、数据搬运机制、原生数值格式和可编程接口可能差异很大。名称相同的 INT4，也可能对应不同打包方式、缩放粒度或执行路径。

先为项目建立一张约束表：

实际支持范围应写成：模型结构 × 输入形状 × 数值格式 × 芯片 × 软件版本 × 执行特性。某个 NPU 支持矩阵乘，并不足以证明它能高效执行任意模型的注意力、状态更新和完整生成循环。

2.2 Prefill 与 decode 给编译器不同的工作负载

预填充（prefill）处理一段输入；解码（decode）利用已有的键值缓存（key-value cache，KV cache），逐步生成新 token。对常见稠密 Transformer，小批量 decode 往往更受权重/KV 访存和固定开销影响；较长 prefill 则通常提供更大的矩阵计算规模。具体瓶颈仍需测量。

两阶段采用不同执行图或编译策略很常见。2026 年公开的 TPU-MLIR 大模型编译论文进一步区分首次 prefill、携带历史 KV 的 prefill 和 decode，给出了将生成流程映射成多个编译入口的具体例子。

2.3 先估算容量，再讨论优化

对各层配置相同、缓存完整历史 K/V 的稠密 Transformer，可以先做以下近似：

$$ M_{\mathrm{weights}} \approx P\frac{b_w}{8},\qquad M_{\mathrm{KV}} \approx 2LBT H_{kv}d_h\frac{b_{kv}}{8}. $$

其中，$P$ 是参数量，$L$ 是层数，$B$ 是批大小，$T$ 是保存的历史长度，$H_{kv}$ 是 KV 头数，$d_h$ 是头维度，$b_w$、$b_{kv}$ 是位宽；此处假设 K/V 的头维度与位宽相同。分组查询注意力（Grouped-Query Attention，GQA）应代入 KV 头数，而非查询头数。多头潜在注意力（Multi-head Latent Attention，MLA）、滑动窗口注意力，以及混合注意力与状态空间模块的模型，需要按实际保存的状态另外计算。

例如，假设一个模型有 30 亿参数，权重全部以 4 bit 保存，原始权重约为 1.5 GB，约 1.40 GiB。另假设 28 层、8 个 KV 头、头维度 128、单请求、8192 token 历史、FP16 KV，则 KV 约为 896 MiB。这些是假设配置的容量计算，尚未计入量化元数据、未量化层、激活、工作区、重复打包和运行时预留。

更实用的总预算是：

1
2

常驻权重 + 持久状态 + 激活峰值 + 内核工作区
        + 格式转换/重复副本 + 运行时与系统开销

编译器能复用部分临时存储，但不能把仍在使用的 KV 当成普通临时激活回收。模型包的磁盘大小，也不能直接当作加载后的内存需求。

2.4 优化目标要同时包含启动与持续运行

端侧需要至少同时观察：编译时间与主机内存、安装包/产物大小、冷启动、首 token 延迟（Time to First Token，TTFT）、相邻输出 token 的间隔（Inter-Token Latency，ITL）、峰值内存，以及设备热稳定后的持续性能。TTFT 应注明从请求到达、模型调用还是其他时点开始计时。

TOPS（每秒万亿次运算，tera operations per second）是运算速率单位；厂商公布的峰值还取决于数值格式、稀疏性和运算计数口径。若 decode 的主要时间消耗在权重读取或图边界上，进一步提高矩阵乘峰值算力可能作用有限。架构选择应依据约定质量与资源限制下的完整生成成本，并把编译、模型更新和多设备维护计入其中。

三、把编译栈分层：表示、优化、执行各负其责

flowchart TB
    A[模型与参考结果
结构 · 权重 · 生成语义] --> B[捕获与导出
图 · 类型 · 形状约束]
    B --> C[图优化与分区
分解 · 量化 · 融合]
    C --> D[可扩展编译后端
张量 IR · 分块 · 内存规划]
    C --> E[厂商编译接口
SDK 图 · 图上下文]
    D --> F[编译产物
目标代码或设备命令
序列化图上下文]
    E --> F
    F --> G[运行时与驱动
加载 · 状态 · 缓冲区 · 同步]
    G --> H[CPU / GPU / NPU 协同执行]

3.1 IR 保存的是程序信息，层级取决于要做的决策

中间表示（Intermediate Representation，IR） 是编译过程中用来表达和变换程序的数据结构。高层表示保留模型算子与张量关系，较低层表示逐渐明确循环、内存和设备操作。Lowering 指向更低层表示的转换；本文保留英文，并用“逐层转换”解释其含义。它本身不表示降低数值精度。

这些对象并不构成必须逐个经过的统一流水线。其中 ATen 提供张量算子，FX 提供图表示与变换工具，导出图可以使用 ATen 算子集。PyTorch Export IR 规范

LLVM IR 也不等同于目标指令集，但可携带目标的数据布局、调用约定和内建函数（intrinsic）等信息。LLVM IR 语言参考

StableHLO、TOSA 与 ONNX 有不同的语义和生态定位；一个厂商后端也可能直接接收子图，内部完成剩余编译。StableHLO 的定位和 IREE 的编译阶段提供了两个不同层面的观察入口。

3.2 理解 MLIR，先理解少量结构

MLIR 的**方言（dialect）**组织相关操作、类型和属性的定义。**操作（operation）**可以表示算术、函数、内存访问等不同语义；**值（value）**由操作结果或块参数产生；**块（block）和区域（region）**组织程序结构，操作还可以包含区域。

在**静态单赋值形式（Static Single Assignment，SSA）**中，每个 SSA 值只定义一次，便于追踪定义与使用关系。但一个值可以引用可变缓冲区，SSA 并不禁止修改其中的内存内容。MLIR 语言参考

Pass 是编译器中的一个处理阶段，可执行分析或变换；多个 pass 组成编译流水线（pass pipeline），本文保留英文 pass。编写 pass 时不仅需要找到某种图结构模式，还要检查类型、属性、副作用和使用关系。MLIR 的接口（interface）、描述可复用性质与约束的 trait、验证器（verifier）和模式重写（pattern rewriting）机制帮助表达这些条件。MLIR Toy 教程适合从小程序逐步建立这些概念。

3.3 导出、分区与代码生成是不同扩展点

• 前端导入解决“模型怎样进入编译器”，需要正确表达模型的算子、常量、形状和状态。
• **图分区（graph partitioning）**解决“哪些操作组成子图、交给哪个后端”，还要考虑精度、形状、依赖与边界成本。
• 代码生成解决“这段计算最终怎样执行”，可以生成内核，也可以调用已有库或厂商编译器。
• 运行时接入解决“产物怎样被加载和调用”，涉及缓冲区、生命周期、同步和错误传播。

例如，ExecuTorch 的分区器（partitioner）标记待委托给后端的图节点，后端处理编译信息，运行时再执行相应产物；这些责任可以分别扩展。ExecuTorch 后端与分区设计

四、发展脉络：从算子优化走向有状态的大模型程序

下面选择公开论文、产品能力和框架发布作为时间节点；年份表示该条事件的公开时间，不表示相关思想第一次出现。不同路线长期并行发展。

从这些节点可以看到一条持续发展的脉络：算子实现 → 跨层表示与优化 → 硬件后端协作 → 有状态程序与交付接口。新阶段并没有消除旧问题；矩阵乘、布局和内存仍然重要，但它们需要放回完整生成流程里评估。

五、前端与形状：先把模型语义完整带进编译器

5.1 torch.compile、torch.export 与 ONNX 的职责

torch.export 会记录保证导出图成立所需的形状约束；若声明的动态范围与程序推导出的约束冲突，导出可能失败。基于一组示例输入成功导出，不代表导出图适用于任意输入；还要遵守形状范围、被特化的参数及其他导出约束。torch.export API

工程上先导出一层或一个小模块，更容易定位问题。完整 generate() 还包含采样、停止条件和循环；有的系统保留这些控制在主机侧，有的将一部分下沉，不能从“模型已导出”推断控制逻辑的位置。

5.2 动态形状是一组决策，不只是一个开关

例如把输入长度分成 128/512/2048 三个形状桶时，一个实际长度为 129 的输入可能填充到 512 桶。填充占比容易计算，但耗时比例还取决于后端是否跳过无效区域、注意力实现和 KV 访问范围。应分别统计逻辑有效长度、实际执行形状和预留容量。

“接受可变长度输入”也不等于“所有底层内核都按完全动态形状执行”。OpenVINO GenAI 的 NPU 文档同时讨论静态形状策略、上下文容量配置及新增的动态输入支持，正好说明这些概念需要分层理解。

5.3 把状态接口写清楚

对一个带 KV 缓存的完整语言模型前向入口，至少要明确下列信息；它们可以作为显式参数，也可以由运行时的状态对象管理：

1
2
3

输入：本次 token ID（或嵌入）、已有 KV、有效历史长度、位置索引、掩码
输出：词表上的 logits、更新后的 KV 或状态更新约定
约束：本次长度 + 历史长度 ≤ 容量；会话与缓冲区归属明确

这里的 logits 是语言模型输出头（LM head）产生的未归一化词表分数。单独的注意力模块通常接收隐藏状态或 Q/K/V 张量，输出注意力计算结果；验证它时应比较相同位置的输出张量。只有接回模型其余层与输出头，才应比较词表 logits。编译器也可能只保留所需位置的 logits，接口应注明输出范围。

图上的“返回更新后 KV”，既可能实现为一次完整复制，也可能通过别名（aliasing）和原地更新（in-place update）复用存储。前端必须正确表达语义，后端和运行时才有机会选择低成本实现。过早抹掉写入关系，或误把状态当成常量，都会使后续优化失去依据。

5.4 控制流、数值与算子分解需要一起检查

常见接入问题包括：依赖张量值的 Python 分支、不可导出的自定义操作、旋转位置编码（Rotary Position Embedding，RoPE）的位置处理、掩码广播、GQA 中查询头与 KV 头的映射、归一化精度以及采样配置。GQA 的 KV 共享不一定需要把数据物理复制到每个查询头，应检查后端的实现方式。对每个问题，先定位是表达能力不足，还是后端没有覆盖对应操作。

算子分解可以把高层操作改写成通用原语，也可能破坏后端对专用融合算子的识别。因此，应先知道后端接受什么，再决定在哪一层分解。表达得更细，不一定更容易优化。

六、图优化与量化：变换必须保持哪些约定

6.1 优化先检查合法性，再估计收益

常量折叠（constant folding）、公共子表达式消除（Common Subexpression Elimination，CSE）、死代码消除（Dead Code Elimination，DCE）、算子分解与融合，都是常见图变换。但一次变换是否成立，取决于类型、形状、数值规则、副作用及结果的使用方式；合法之后，还需单独评估收益。

以“融合矩阵乘和残差相加”为例，数学表达可以合并，但融合实现可能改变中间舍入精度。另一个消费者若仍需要矩阵乘的原始结果，也会影响能否完全消除中间写回。优化收益需要针对完整子图测量。

6.2 模式重写与合法化各自解决什么

模式重写（pattern rewriting） 匹配一段程序并替换它；合法化（legalization） 将程序转换到目标允许的操作和类型集合。目标集合可以同时包含多个 dialect，合法也不等于已经生成硬件代码。

一个可维护的变换应记录四件事：匹配条件、拒绝条件、替换结果和验证方法。比如对连续转置，只在置换互逆、张量语义允许且相关属性一致时消除；对不满足条件的案例，应该保持原图或明确报错。

MLIR 的 Dialect Conversion组织转换目标、重写规则与类型转换。学习时要特别区分“规则没有匹配”“流水线仍有非法操作”和“转换后的程序数值错误”，这三种问题的排查入口不同。

以 Toy 教程的二维张量转置为例，下面的片段保留了显式类型，便于观察 SSA 使用关系：

1
2
3
4

// 变换前：同一二维张量连续转置两次。
%0 = "toy.transpose"(%input) : (tensor<2x3xf64>) -> tensor<3x2xf64>
%1 = "toy.transpose"(%0) : (tensor<3x2xf64>) -> tensor<2x3xf64>
func.return %1 : tensor<2x3xf64>

1
2

// 变换后：返回原输入，结果的类型与数值语义保持一致。
func.return %input : tensor<2x3xf64>

这里利用的是 Toy 操作的纯张量语义。若 %0 还有其他使用者，第一条转置仍需保留；若换成带任意排列属性的高维转置，则必须检查两次排列是否互逆。张量层的等价改写，接入有可变别名、布局或状态副作用的更低层表示时，也需要重新审视前提。Toy 的模式重写示例

6.3 数值格式怎样进入 IR 与执行约定

编译器需要区分数值含义、逻辑类型、物理存储与计算类型。FP16、BF16 都是 16 位浮点，分配给指数和尾数小数部分（fraction field）的位数不同；INT8、INT4 是整数表示，量化后还需由缩放因子（scale）、零点（zero-point）等参数说明它们代表什么数值。基础格式对比与误差算例见推理与部署篇第 6.2 节。本节继续追踪这些信息怎样变成可执行程序。

均匀仿射量化的关系是：

$$ q=\operatorname{clip}\left(\operatorname{round}(x/s)+z,q_{\min},q_{\max}\right), \qquad \hat{x}=s(q-z),\quad s>0. $$

其中 $q$ 是存储整数，$\hat{x}$ 是它表达的近似实数。round 表示按约定规则舍入，clip 表示将超出量化整数范围的值限制到边界；这里的范围裁剪（clipping）对应饱和处理（saturation），不同于向零截断（truncation）。

MLIR 的 i4 是不带有符号/无符号标记的 4 位整数类型（signless integer type），符号解释取决于操作或量化类型的约定，不能单独表达完整量化语义。MLIR 整数类型

MLIR 的量化类型区分存储类型（storage type）与表达类型（expressed type）：前者保存量化编码，后者表示这些编码对应的数值所采用的类型，通常为浮点。表达类型并不自动决定硬件累加类型。MLIR 量化设计

调试时可以先把约定打印成与具体 API 无关的清单：

1
2
3
4
5
6

storage_type: signed INT4       expressed_type: FP16
code_range: [-8, 7]             group_axis: K
group_size: 128                 scale_type: FP16
zero_point: 0                  rounding: nearest, ties-to-even
packing/layout: 由产物和后端接口共同约定
compute/accumulator/output: 由已选择的执行路径明确记录

这是一份跨层核对清单，不是可以直接解析的 MLIR 类型声明，也不代表任意 NPU 都支持这些属性组合。FP16/BF16 转换同样要保持类型语义：二者位宽相同，位重解释（bitcast）通常会改变数值，不能替代数值转换；涉及累加或量化缩放时，更需要核对数值转换和存储重排的区别。

flowchart TB
    A[模型数值与量化约定
scale · 零点 · 轴 · 分组] --> B[IR 类型与 Q/DQ
形状 · 转换 · 合法化]
    B --> C[物理打包与布局
INT4 字节 · 元数据 · 对齐]
    C --> D[内核与设备执行
读取 · 转换 · 乘法 · 累加]
    D --> E[输出与状态写回
输出类型 · 重量化 · KV]

6.4 INT4 怎样装进字节，怎样从内存中读回来

以有符号补码 INT4 为例，1100 表示 −4，0010 表示 2。ONNX 的紧凑存储规则是：连续两个元素中，第一个放在字节的低 4 位，第二个放在高 4 位；元素数为奇数时，最后补半个字节。厂商内部布局可能另外重排，不能把导出文件中的顺序直接当作设备格式。ONNX INT4 存储规范

假设权重 [-1.0, 0.0, 0.5, 1.5] 用 s=0.25, z=0 表示，得到整数 [-4, 0, 2, 6]：

1
2
3
4
5

逻辑顺序：      -4        0        2        6
4 位补码：    1100     0000     0010     0110

第一个字节：  0000 1100 = 0x0C   （低 4 位放 -4，高 4 位放 0）
第二个字节：  0110 0010 = 0x62   （低 4 位放  2，高 4 位放 6）

下面的小实验只依赖 Python，用于检查位序和符号恢复：

1
2
3
4
5
6
7
8
9

q = [-4, 0, 2, 6]  # 本例固定为四个有效 INT4 值
packed = bytes((q[i] & 0xF) | ((q[i + 1] & 0xF) << 4)
               for i in range(0, len(q), 2))
codes = [(byte >> shift) & 0xF for byte in packed for shift in (0, 4)]
restored = [code - 16 if code >= 8 else code for code in codes]
dequantized = [0.25 * value for value in restored]
assert packed.hex() == "0c62"
assert restored == q
assert dequantized == [-1.0, 0.0, 0.5, 1.5]

若把取出的 1100 当无符号整数 12，再乘 0.25，就会得到 3.0，而非 −1.0。这个错误来自符号解释；若把高低半字节颠倒，则是元素顺序错误。两者都不是量化算法本身的误差。

逻辑张量形状与物理缓冲区大小需要分开记录。 连续打包的 N 个 INT4 只需 ceil(N/2) 字节保存编码，但仍要另外保存或约定 scale、零点、组大小和有效元素数。比如 128 个 INT4 加一个 FP16 scale，共为 64 + 2 = 66 字节；若某个假设布局要求每组独立按 16 字节对齐，则这组会占 80 字节。实际是否按组、按行或按整个张量对齐，由格式决定。

因此，布局变换必须同时处理权重编码和元数据。当原来沿 K 维的分组被转置、分块或交错存储时，scale 仍需对应原来的那组元素。把 4 位值装进 uint8/int32 容器，也不会自动改变它们的逻辑数值类型。

6.5 分组量化、累加类型与矩阵乘怎样配合

先明确计算路径。 W4A16 表示 4 位权重、16 位激活，还需说明具体编码及 FP16/BF16 的选择。一种实现是在内核中按块把权重反量化为浮点，再与浮点激活计算；另一类量化路径会让权重和激活都进入受支持的整数或低比特矩阵运算。不能只根据磁盘上的 INT4，推断乘法、累加或 KV 的类型。

下面以激活和权重都采用均匀整数量化为例。对一个分组点积，若组 g 的参数为 $s_{a,g},z_{a,g},s_{w,g},z_{w,g}$，则：

$$ y\approx\sum_g s_{a,g}s_{w,g} \left[\sum_{k\in g}(q_{a,k}-z_{a,g})(q_{w,k}-z_{w,g})\right]. $$

不同组的整数累加结果，需要按各自尺度组合。做一个简化算例：激活尺度为 1、零点全为 0，两组整数点积结果都为 4，权重尺度分别为 0.5、0.25；结果应是 4×0.5 + 4×0.25 = 3。若先合并成 8，再统一乘 0.5，就变成 4。这个差异说明分组边界是计算语义的一部分。

编译器需要确认：

1. 分块的 K 维与量化分组是否对齐；跨组的 tile 怎样读取多个 scale。
2. 零点修正怎样实现；若被折叠进其他项，是否仍保持同一数学关系。
3. 整数部分和用什么类型累加，再在哪里缩放、转换或合并。
4. 内核需要的权重布局与元数据布局是否相容，尾块怎样处理。

累加器需要单独选择。 例如 4096 项 127×127 的和为 66,064,384，超过 INT16 范围，因此 INT8 输入不能推出 INT8 或 INT16 累加。INT32、FP32 等较宽类型可服务相应路径，但仍需检查归约长度、数值范围和目标指令；浮点累加更宽也无法恢复输入量化时已经丢失的信息。

AWQ、GPTQ 等方法帮助选择权重量化参数与降低误差；要落到设备，仍须匹配分组、元数据、内核与图变换。AWQ、GPTQ

6.6 Cast、Q/DQ、重量化与融合的边界

Q/DQ 分别指量化与反量化（Quantize/Dequantize）。它们与普通类型转换、位重解释和内存重排有不同语义：

对上一节的整数 −4，普通 INT4→FP16 数值转换得到 −4.0；使用 s=0.25, z=0 做反量化才得到 −1.0。编译器不能把这两种操作互换。MLIR 量化类型转换

再看最简单的逐张量整数乘加：忽略偏置（bias），设已完成零点修正的累加值为 acc，输出需要采用尺度 $s_y$、零点 $z_y$，则理想重量化关系为：

$$ q_y=\operatorname{clip}\left( \operatorname{round}\left(\frac{s_a s_w}{s_y}\operatorname{acc}\right)+z_y, q_{y,\min},q_{y,\max}\right). $$

目标实现可能用定点乘法与移位近似这一过程，因此需要检查乘法中间位宽、舍入和饱和。以整数量化为例，ONNX QuantizeLinear 采用舍入到最近值、中点取偶数（round to nearest, ties to even）；它与直接向零截断不是同一规则，例如 1.9 分别得到 2 和 1，2.5 的中点取偶结果为 2。超范围后饱和到边界，也与整数回绕（wraparound）不同。ONNX QuantizeLinear 的舍入与饱和约定

Q/DQ 可以作为图上的数值边界，后端再把解包、缩放、矩阵运算和输出转换融合在一个内核中。融合可以消除独立临时缓冲区，但应保留约定的数值行为；直接删除一个有损的 Q→DQ 往返，会改变程序。若后端缺少匹配实现，则可能拒绝、回退到其他执行路径（fallback），或先展开为浮点再执行。TensorRT 显式量化与 Q/DQ 融合

因此，检查转换后的 IR 还不够。还需确认实际执行类型、转换位置、是否生成全量展开副本，以及融合是否改善完整子图的时间与内存。

6.7 从量化数值验证到编译产物验证

先保存三组结果：原始参考计算、采用相同量化参数的数值参考、编译后执行。第二组与第一组的差异帮助评估量化本身；第三组与第二组的差异帮助定位编译或后端问题，还要考虑已声明的计算精度与归约顺序。

只校准短输入，未必能代表长上下文、图像 token 或多轮历史下的分布。校准与质量评测集应分开，按阶段检查误差，再决定哪些层、哪些状态需要保留较高精度。KV 量化还需要独立确认量化粒度、静态或动态缩放、布局与读写约定，不能从权重量化配置直接推断。

这一部分的 P0 是格式语义、参数传递、打包/布局、转换与累加的区别，以及逐阶段验证；P1 是所选后端的具体实现；其他厂商编码和特殊低比特格式按需深入。第十三节第三周先用小实验建立这些约定，第五、六周再把它们接入布局与真实后端。

七、张量程序、布局与内存：把计算映射到硬件

7.1 从 MatMul 到分块执行

矩阵乘 $C_{M\times N}=A_{M\times K}B_{K\times N}$ 的计算量约为 $2MNK$ 次算术运算，一次乘加计作两次。硬件通常不能把完整矩阵都放进片上存储，因此需要计算分块（tiling），使读入的数据在被换出前获得足够复用。

一种分块策略至少决定：输出块大小、归约块大小、遍历顺序、数据进入哪级存储、多个执行单元怎样分工，以及何时同步。TensorIR 将块、迭代与张量计算原语组织成可变换的程序；MLIR 的 Linalg 等表示也可以保存结构化计算信息。TensorIR 论文

这里的**计算调度（schedule）**指张量程序的执行组织方式，包括循环顺序、分块、并行和数据搬运等安排，与推理服务的请求调度含义不同。张量化（tensorization） 通常指把一段匹配的计算映射到硬件提供的张量/矩阵计算原语，也不同于将 Python 列表转换为 Tensor。

7.2 一次片上容量计算

设一个输出分块为 $M_t=32,N_t=64$，归约分块 $K_t=128$，A/B 使用 2 字节元素，累加器使用 4 字节：

1
2
3
4
5
6

A 块：32 × 128 × 2 =  8 KiB
B 块：128 × 64 × 2 = 16 KiB
C 累加器：32 × 64 × 4 = 8 KiB
合计：32 KiB

A/B 都采用双缓冲：2 × (8 + 16) + 8 = 56 KiB

若假设 A/B 块与 C 累加缓冲区都占用同一个 48 KiB 的片上存储池，前一种组织在容量上可能可行，双缓冲（double buffering）则已经超限，尚未计入对齐、scale 和其他临时量。需要缩小分块、改变存储层级或调整流水线，不能只看“开启双缓冲可能隐藏搬运”这一收益。若累加器位于独立寄存器文件或专用存储中，应分别核算资源，不能机械地相加。

这个算例只说明资源预算，不指定某款 NPU 的原生精度或存储分配方式。实际分块还可能受矩阵单元形状、存储体（memory bank）、直接内存访问（Direct Memory Access，DMA）的传输与对齐要求，以及编译器约束影响。

7.3 布局是跨算子的约定

数据布局（data layout） 描述逻辑索引如何映射到物理存储。一个 MatMul 希望的权重块布局，可能与通用导出格式不同；注意力计算对 KV 的访问方向，又可能与追加写入最方便的方向不同。

局部布局优化要核算整条数据路径。若某个内核快了，但每一层都新增转置、打包或 CPU/NPU 往返，端到端可能退化。对低比特权重，还要确认“转置数据”之后 scale 的轴和分组语义是否同步改变。

有些 reshape/transpose 在某种表示里只是视图，在后端不接受相同步长或打包方式时却需要实际复制。分析时要查看执行后的访存和转换，不能仅凭高层算子名称估计成本。

7.4 Bufferization、内存规划与状态管理

这三个概念彼此关联，但负责不同的问题：

MLIR 的 One-Shot Bufferize围绕别名和读写冲突决定原地（in-place）或非原地（out-of-place）处理；ExecuTorch 的内存规划根据张量大小与存活期安排预分配内存区（arena）内的位置。嵌入的厂商子图仍可能管理自己的内部内存，外层规划不等于掌握了全部设备占用。

这里将“存活期”用于描述值仍可能被后续计算使用的范围；活跃性分析（liveness analysis）帮助判断各程序点哪些值仍然需要保留。“缓冲区生命周期（buffer lifetime）”则涉及分配、使用到释放的完整过程，异步设备的实际使用也必须计入。

例如，一份输入被两个并行任务借用，第二个任务尚未完成时，即使主机代码已提交了后续操作，也不能立即复用该区域。存储复用需要同时满足数据依赖和执行完成条件。

7.5 融合、层组和流水线需要共同优化

算子融合（operator fusion）可以减少中间结果写回和任务提交；TPU-MLIR 等工具中的层组（layer group）将一组操作组织起来，结合切片与局部存储规划复用数据，不一定意味着合成一个内核；流水线（pipelining）则尝试重叠搬运和计算。它们都会影响片上资源需求与并行程度，具体语义要按工具实现理解。

“融合越多越好”并不成立。更大的融合区域可能超过片上容量，迫使中间数据反复写回片外内存，或采用重计算（recomputation）；寄存器压力过高也可能引发寄存器溢出到内存（register spilling），这些都与数值溢出（numerical overflow）不同。某个小内核单独执行反而更灵活。TPU-MLIR 的官方仓库视频索引中，LayerGroup 教学适合观察计算组织与局部内存之间的关系。

自动调优也应围绕这一成本模型。搜索空间可以包含分块、顺序、布局和并行参数，代价由实机测量或预测模型给出；测量要控制频率、温度和输入条件。较早的 Ansor展示了程序生成与搜索的思路，学习概念后再对照所用版本的调优接口。

八、工具链地图：哪些需要深学，哪些先理解定位

8.1 先分清基础设施、执行栈和专用编译器

MLIR 提供构建编译器的基础，TVM、IREE 与专用编译器则选择自己的程序组织和执行方式。ExecuTorch、LiteRT、ORT 主要承担不同的模型与执行接入路径，可能调用厂商编译器。它们不能只用一张速度排行榜比较。

8.2 MLIR 与 TVM 怎样选

若目标是理解 NPU 编译器实现，本文建议以 MLIR 基础 + 一个可阅读的目标编译器为主线，再用 MLC 课程中的 TVM 示例建立张量程序直觉。若已有项目使用 TVM，则可以反过来，以项目栈为主，再学习 MLIR 中相同问题的表达方式。

当前 TVM 文档以 Relax 与张量 IR 的协作为重要路径，并出现 TIRx 等较新的接口；旧课程中的 Relay、tvm.tir 或早期 TVMScript 代码不能假定直接适配新版本。先固定课程环境，或明确进行版本迁移。TVM 架构、Relax

8.3 BYOC、delegate 与 EP 是可比较的接入思路

TVM 的 Bring Your Own Codegen（BYOC，外部代码生成后端接入）、ExecuTorch 的**后端委托执行（delegate）**和 ORT 的 EP 都允许外部后端承担一部分程序执行，但它们的 IR、生命周期和接口不同。

学习时可以用同一组问题比较：怎样声明能力，怎样选择子图，怎样序列化产物，谁分配输入输出，谁处理同步，失败后怎样恢复。共同的概念可以迁移，接口实现需要按项目重做。

一个细节很能说明验证的重要性：TVM 当前 BYOC 教程中的示例 NPU 后端是教学用桩实现（stub），只记录调度决策，不执行实际计算，输出缓冲区也未初始化，因此该部分只检查形状。跑通教学后端，证明的是接口链路；数值正确性和 NPU 性能需要另外验证。

8.4 ONNX、StableHLO 与 TOSA 要学到什么程度

它们首先是语义与接口边界。需要了解支持哪些算子、类型、形状与量化方式，以及生产者和消费者的版本约定；只有编写导入器、转换器或后端时，才需要深入更多算子规范。

StableHLO 的兼容保证针对其规定的可移植产物和语义范围，不能外推为设备二进制兼容。ONNX 的算子版本也不代表某个 EP 必然实现对应操作。TOSA 可以作为专用加速器的算子接口选择，最终仍依赖具体后端。StableHLO 兼容说明

8.5 同名与相邻技术的边界

• MLC 可以指机器学习编译这一领域、MLC.ai 课程，或 MLC LLM 项目；阅读时明确上下文。
• TPU-MLIR 在本文指 Sophgo 的开源项目，不是 Google TPU 编译器。
• Triton language、TileLang 等 DSL 可以帮助理解计算分块和内核实现；是否能用于目标 NPU，要看具体后端支持。
• CUDA/CUTLASS 的优化经验可以帮助建立 GPU 性能直觉，移植到 NPU 仍需重新理解指令、存储和同步。
• 运行时支持某芯片与开放该芯片的编译器或内核开发接口是不同能力，需要分别确认。

对同类框架，先深入一套，再比较另一套怎样表达同一个问题；这样才能识别设计取舍，而不只是累积 API 名称。

九、端侧 NPU 生态：厂商工具链怎样连接模型与设备

9.1 按设备与开放层次比较

下表列出可进一步验证的路线，表示项目定位与公开入口，不保证任意模型、量化格式和芯片组合都能运行。

对应资料入口：Qualcomm 后端、MediaTek 集成、RKLLM、TPU-MLIR、OpenVINO NPU、Ryzen AI OGA、Core ML 状态模型。

9.2 Qualcomm：接入层可以不同，厂商编译边界仍需理解

ExecuTorch、LiteRT 和 ORT 为 Qualcomm 设备提供了不同的接入方式。它们在模型表示、分区和运行时接口上不同，底层仍需配合目标 SDK 与设备能力。

ORT QNN EP 的公开文档给出了一条很有学习价值的路径：图经过处理和编译后，可以缓存 QNN 上下文二进制文件（context binary），以减少后续会话创建成本；还提供关闭 CPU EP 回退的配置，帮助验证图是否完整由指定 QNN 后端承接。QNN 还可选择不同执行后端，HTP（Hexagon Tensor Processor）路径与 CPU/GPU 路径应分开核对；关闭 ORT 层的 CPU 回退也不等于得到了设备内部逐指令执行轨迹。QNN EP 文档

这说明后端接入至少包含两类工作：把模型表达成可接受的图，以及管理编译后上下文的加载与生命周期。低层编译器未必全部可修改时，图改写、量化、分区、缓冲区和负载配置仍然是重要优化位置。

9.3 Rockchip 与 Sophgo：交付接口与编译源码的不同观察点

RKLLM 的公开工作流是 PC 侧转换/量化、板侧通过 Runtime 执行。其模型文件、转换选项和 C/C++ 接口适合研究实际部署约束；通用 RKNN 与 RKLLM 应按任务选择，不能因为都面向同一品牌芯片就互换流程。RKLLM 官方说明

TPU-MLIR 则提供面向目标芯片的编译源码和开发文档，可以继续追踪从高层图到目标表示的变换。做模型部署时关注转换与运行；研究编译实现时，可选一个操作沿导入、lowering、层组和代码生成路径阅读。具体底层组件的开放范围仍需逐项检查。TPU-MLIR 开发入口

两条路线对应不同的学习抓手。已有 Rockchip 设备，可以先把状态、精度与性能测清楚；希望修改实际编译 pass，又没有既定厂商约束时，可以从具有完整示例和测试入口的编译项目着手。

9.4 AI PC：上层部署与底层可编程能力分开看

Intel NPU Compiler 提供编译器源码、MLIR 入门、构建和调试指南，使学习者有机会从 OpenVINO 的模型使用进一步进入编译实现。驱动、固件、模型执行层与编译器仍有各自的版本和接口。Intel NPU Compiler

AMD Ryzen AI 的官方 OGA 文档区分 NPU-only 与 NPU/iGPU 混合模式；其可用模型与处理器组合需要查支持范围。底层的 MLIR-AIE/IRON则提供面向 AI Engine 的编程和编译入口。能够运行一个预优化模型，与能够自行编写该设备上的计算程序，应分别验证。Ryzen AI 执行模式

9.5 Apple：状态接口值得学习，执行位置需要证据

Core ML 的有状态模型（stateful model）能表达 KV 等跨调用状态，适合学习状态归属、模型接口和反复调用的成本。但其公开状态示例包含 CPU/GPU 配置，不能把该示例的收益直接解释为 Apple Neural Engine（ANE）收益。

配置允许的计算单元，也不是每个操作实际运行位置的证明。应使用模型分析与性能工具检查设备分配，并在同一模型、输入和精度下比较。Core ML 状态示例、WWDC24 的状态模型与性能工具讲解

9.6 当前发展值得关注的三件事

从上述公开资料中，可以提取三个方向：第一，导出与厂商后端的接口越来越完整；第二，KV 和生成阶段成为明确的编译/运行时对象；第三，部分专用 NPU 编译器提供了更深入的源码与开发入口。 这是对这些项目的归纳，不意味着所有芯片的软件栈都已具备相同能力。

新的统一 API 会降低接入成本，但形状、布局、数值格式和设备内存的差异仍然存在。学习时应追踪这些约束怎样被表达和传递，而不是只跟踪 SDK 名称变化。

十、编译与运行时的结合点：状态、边界和产物

10.1 多个编译入口共享同一份逻辑状态

flowchart TB
    P[首次输入
Prefill 入口] --> K[会话 KV 状态
布局 · 容量 · 有效长度]
    K --> D[单步 Decode 入口]
    D --> K
    K --> H[带历史的 Prefill 入口
新一轮输入或后续输入块]
    H --> K
    P -- 提示词已处理完成 --> O[所需位置的 Logits
选择下一 token · 停止判断]
    H -- 提示词已处理完成 --> O
    D --> O
    O -- 继续生成：下一 token --> D
    O -- 停止或取消 --> R[等待执行完成
释放或保留会话状态]

常见自回归生成中，完整提示词的 prefill 结束后，利用最后一个有效位置的 logits 选择第一个输出 token。后续 decode 将上一步选出的 token 输入模型、更新 KV，再产生用于选择下一个 token 的 logits。因此，刚选出的 token 通常要到下一次前向计算才进入 KV。分块 prefill 在提示词处理完成前，只接续计算与状态，不开始生成回复。

这个接口有几个不能遗漏的条件：prefill 与 decode 使用相容的 KV 表示；有效长度与写入位置一致；容量超限行为明确；不同会话的写入互不污染；执行未结束时不能提前回收缓冲区。前缀缓存复用可以共享只读 KV 块，但共享后再写入时，需要独立缓冲区或写时复制（copy-on-write）等隔离机制，不能直接覆盖其他会话仍在使用的数据。PagedAttention 的 KV 共享与写时复制

固定容量的 KV 并不意味着注意力可以读取全部区域。尚未写入的位置需要通过正确的长度和掩码排除。同一模型跨形状桶切换时，要确认是继续复用、转换状态，还是重新计算；切换到不同模型或权重时，通常应重建 KV，不能仅凭形状相同复用。

10.2 原地更新、共享内存与零拷贝

“图中使用同一个 KV 参数”“主机与设备共享物理内存”“调用之间没有全量复制”是不同层面的事实。零拷贝（zero-copy）需要满足可访问性、对齐、地址注册、布局和生命周期等条件；即使共享物理内存，仍可能有缓存一致性维护（cache coherency maintenance）、同步和格式转换。

ExecuTorch 的 Qualcomm 示例提供了共享缓冲区接入说明，可以追踪缓冲区申请、注册、使用与释放的完整过程。Qualcomm 示例中的共享内存流程

函数返回不一定表示设备任务已完成。 异步执行时，状态和输入输出缓冲区必须保持有效，直到相应的设备任务完成；取消请求通常也需要处理已经提交的设备工作。编译器的存活期分析和运行时的完成条件，应在这一边界对齐。

10.3 子图覆盖率必须换算成边界成本

flowchart TB
    A[NPU 子图 A] --> B[等待设备完成
转换布局或准备缓冲区]
    B --> C[CPU 上的未覆盖操作]
    C --> D[准备设备输入
提交与同步]
    D --> E[NPU 子图 B]
    E --> F[后续模型层与解码步
边界可能反复出现]

一个只占少量计算的操作，可能切断较大的融合区域，或在每层每步引发切换。因此，按“支持算子个数”或“图节点覆盖率”衡量适配进度，可能低估最重要的问题。

可以用以下估算初步排序：

$$ T_{\mathrm{boundary/token}}\approx L n_b t_b. $$

假设 28 层，每层两次边界，每次平均 50 微秒，则每个 decode token 增加约 2.8 毫秒，128 个 decode 步约增加 358.4 毫秒。这是串行、均匀边界的分析算例；实际需要区分可重叠操作与关键路径，也要测量转换和同步的真实成本。

实际迁移收益应估为“原路径时间减去新设备计算、转换、搬运和额外同步时间”。有时候补一个小操作的覆盖，比把已有大算子再优化一点更有价值；有时候将一组操作整体留在 CPU/GPU 上更合理。

10.4 编译产物是有条件的执行程序

产物清单至少应记录：

1
2
3
4
5
6

源模型与权重校验、分词器与模板
导出器、IR/opset、编译器与优化流水线版本
目标芯片、精度、布局、输入范围和 KV 协议
运行时/驱动/SDK 依赖与已验证组合
权重、子图、外部常量、设备二进制及其对应关系
正确性结果、性能基线、冷启动和回退条件

AOT 产物可能减少首次启动工作，也可能增加按芯片、形状和精度维护的变体数。若针对 4 类目标、3 个形状桶、2 种精度生成独立变体，理论上就有 24 个组合；实际可通过共享权重、按需生成和减少支持范围控制规模，但要验证实现支持。

可移植图和目标二进制通常有不同的兼容边界。编译缓存的标识应包含会影响生成结果的输入，版本变化后根据兼容约定与回归结果决定复用，不能仅用“模型名称相同”作为依据。

10.5 Prefill/decode 分图与跨设备分工要分开决策

两个阶段有不同编译入口，仍可以在同一个 NPU 上执行；跨设备分工则额外涉及 KV 可访问性、布局转换、资源争用与同步。应先分别测清两个阶段，再判断迁移收益是否覆盖状态交换成本。

端侧的 CPU/GPU/NPU 协同还要考虑应用其他负载。例如视觉编码器、语音处理或 UI 可能正在使用 GPU。最短的独立模型延迟，并不一定对应完整应用最稳定的资源安排。

十一、架构选型：围绕目标、边界与长期成本

11.1 三类目标，对应不同起点

如果当前问题发生在导出和 KV 接口，先修改这一层；若热点在厂商编译结果内部，需要判断已有配置、图改写或算子扩展能否解决，再决定是否扩大自研范围。架构选择要和团队实际能维护的层次对应。

11.2 一个决策实例：模型可加载，但长输入慢

先把现象拆开：慢的是首次编译、prefill、KV 初始化、图切换，还是每步 decode？然后逐项比较：

1. 增加一个更合适的输入形状桶，是否减少了填充计算。
2. 使用分块 prefill，是否降低峰值内存，同时增加了调用与状态开销。
3. 将一个小操作纳入后端，是否减少了反复切换。
4. 改变 KV 格式或访问范围，是否改善访存且保持长上下文质量。

每次只改变一个主要因素，并保留失败样例。把所有开关一起调整，即使得到了更好的最终数字，也很难形成可迁移的判断。

11.3 用架构决策记录固定取舍

可以沿用上一篇的架构决策记录（Architecture Decision Record，ADR），增加编译特有的信息：

1
2
3
4
5
6
7
8

问题：哪个模型、输入范围或设备限制触发了决策？
基线：当前图、精度、布局、产物与设备结果。
候选：修改前端 / 修改 pass / 调整后端配置 / 改变设备分工。
正确性：语义约束、量化误差、状态与边界检查。
收益：编译成本、启动、prefill/decode、内存与持续性能。
代价：新增变体、源码维护、调试难度与模型升级成本。
决定：采用哪条路径，在哪些条件下生效？
退出条件：哪些模型/版本变化触发重编译、重测或回退？

一份能说明“为什么只给某组形状启用融合”的记录，往往比“某编译器更先进”的概括更有工程价值。

十二、怎样验证：逐阶段定位，最后回到完整模型

12.1 正确性检查要沿转换链展开

推荐保留以下参照：原始浮点实现、导出图、变换后图、量化数值模型、编译后的执行结果。逐层或逐模块比较，可以找出误差最早出现的位置。

浮点结果未必逐位一致（bitwise identical），应先按数据类型、操作和任务要求确定容差。只比较最终生成文本，很难区分编译错误与采样分歧；逐步比较 logits 时应固定输入 token 序列，避免生成路径分叉掩盖最初误差；只比较短序列的平均误差，也可能遗漏状态和尾块问题。

12.2 Pass 测试与数值测试互相补充

编写图重写时，需要正例证明该匹配被改写，反例证明不符合前提的图不会被误改。FileCheck 一类文本检查适合确认 IR 结构，但不能单独证明数值正确。

随后执行变换前后程序，使用多组输入验证；涉及状态时连续运行多步，涉及形状时覆盖分桶和上下界；最后再接回完整模型。测试应保护变换的语义条件，而不是只确认代码里出现过某个 pass 名称。MLIR Toy 教程中的变换与 lowering

12.3 性能报告分清四个时间窗口

比较前固定模型、数据、生成条件、硬件与版本，分别报告输入实际长度和编译形状。跨精度或跨后端比较时，应补同后端的参考精度基线，避免把引擎差异全部归因于量化。

如果某内核占总时间 20%，把它加速到 2 倍，在其他部分不变时，端到端加速上限约为 $1/(0.8+0.2/2)=1.11$ 倍。算例可以帮助设定预期，但最终仍要检查改动是否影响融合、搬运与状态。

12.4 一份可复查的编译实验记录

1
2
3
4
5
6
7
8

实验编号、源码提交、配置、设备与完整软件版本
模型/权重、导出图、pass 流水线、形状和精度
每阶段 IR 或可取得的诊断产物
量化元数据、KV 表示、分区与回退行为
正确性容差、输入边界、质量结果和失败样本
编译/初始化/prefill/decode 时间与重复测量
峰值内存、布局转换、实际设备轨迹、持续性能
相对基线的唯一变化、收益、退化条件与决定

公开工具有时无法展示固件或内部内核的全部细节。此时应明确可见范围，通过分阶段计时、输入扫描和受控对照缩小原因，不把猜测描述成已经定位的硬件瓶颈。

12.5 把稳定性放进最后一次验收

检查重复加载、会话重置、取消后资源释放、容量超限、设备错误、模型切换和版本回退。异步运行时尤其要覆盖“请求结束，但设备仍在使用缓冲区”的情况。

交付包应带上支持矩阵、编译与运行配置、模型/产物标识和复现说明。能够在已声明的条件下稳定重建和执行，才算完成这条编译工程链。

十三、8 周学习路线：用课程建立理解，用实验贯通编译链

13.1 学习前提与投入边界

这条路线面向已经能使用 Python、PyTorch，并理解矩阵乘、注意力和自回归生成的读者。修改 C++ 编译 pass 还需要基本的 C++ 阅读能力、CMake 构建与调试经验。若这些基础不足，应先补对应内容，或先选择模型接入方向。

安排采用每周 5 个主学习日，每天约 8 小时，共 8 周、约 320 小时。每天只规定任务与产物；第 6、7 天用于休息、补漏和处理构建问题，不叠加新的必修内容。已经完成上一篇推理路线的读者，可以直接复用模型、质量样本和性能脚本，把节省的时间投入 IR、pass 与内存分析。

两个月的目标是：读懂一条编译链，完成一次有条件的程序变换，接入一套目标工具链，并留下能解释正确性、性能与限制的实验。 独立实现完整 LLM 编译器、新芯片后端或全部设备内核，需要更长的积累。

建议采用“带着当天问题看指定章节 → 在固定版本上做小实验 → 对照 IR 与结果解释变化”的顺序。视频负责建立直觉，官方文档负责核对接口，源码和实验负责验证理解。每周至少一半投入留给代码、调试和复盘，避免把完整刷课作为完成标准。

13.2 课程与资料：编译主线优先，模型和 GPU 内容按需补齐

本篇建议的顺序是 MLC 精选 → MLIR 实作 → 一个真实 NPU 工具链。 MIT 量化课穿插在第三周；vLLM 服务短课和 NVIDIA 推理直播可以补部署视角，但在这条编译路线中排在 IR、变换与内存之后。

两条版本说明需要提前记住：MLC 2022 notebook 与当前 TVM API 存在差异，优先保留课程环境复现实验，再对照当前文档理解演进；MLIR 视频用于学习机制，构建教程和源码应取同一个版本。不要为了追上所有仓库的最新提交，把第一周耗在依赖迁移上。

上述视频入口来自课程或项目官方索引；TPU-MLIR 的 B 站入口可能受登录或访问限制，无法播放时使用同主题手册与源码。课程选择以公开教学内容和实验衔接为依据，不要求额外购买付费课程。

13.3 先固定一个实验仓库和一条设备路径

准备两类实验对象，并让它们共享形状、精度、状态和测量约定：

• 小模块：矩阵乘、MLP 或单层注意力。用来观察 IR、写重写规则、检查布局与缓冲区；尺寸小，错误容易定位。
• 小型完整 LLM：选目标后端明确支持、当前设备能容纳的模型。用来验证分阶段生成、量化质量和真实运行边界；模型、权重版本与 tokenizer 在第一周固定。

小模块可以从参考模型中提取；Toy 程序只承担编译机制练习。不要把 Toy 的 CPU lowering 直接描述成完整 LLM 到 NPU 的编译路径。

1
2
3
4
5
6
7
8

compiler-study/
  configs/     # 源码版本、模型、设备、形状、精度与环境
  models/      # 小模块与完整模型的导出入口
  ir/          # 关键阶段 IR、分区与诊断信息
  transforms/  # pass、重写规则或后端适配改动
  tests/       # 数值、形状、状态、重写反例
  runtime/     # 加载、执行、状态与测量脚本
  reports/     # 原始数据、每周结论与最终复现说明

第一周就完成目标 SDK 的可用性检查和最小示例准备。 如果工具包获取、主机环境或设备调试不可行，立即沿 CPU 路径完成共同基础，并把上板实验明确列为待补。模拟器可辅助功能检查，其耗时不能替代真实设备性能。

13.4 八周安排：每天有动作，每周有验收

第 1 周：固定模型语义，导出可检查的程序

本周问题：从 Python 模型到导出图，哪些输入、状态和形状约束必须保持？

本周材料：MLC 中文课程第 1、4 课的概览部分；torch.export 官方教程。注意力基础不足时选看 MIT 第 12 讲；资源核算不足时补 CS336 第 2 讲。回看本文第二、五节。

本周验收：能解释每个输入与输出的含义，并从图中找到一个具体形状约束。性能数据只是后续对照的起点；目标 SDK 暂不可用时，第一周就注明采用 CPU 实验路径。

第 2 周：读懂 MLIR，写一个能证明条件的重写

本周问题：一个图变换为什么合法，编译器通过什么结构表达并检查它？

本周材料：MLIR 入门视频与 Toy 第 2、3、5、6 章。重点是 IR、模式与 lowering；语法解析器沿用教程实现。回看本文第三、六节。

本周验收：能够解释一个规则为什么不应匹配某个反例。构建资源有限时，限定目标和并行度，使用匹配版本的开发环境；本周无需编译 LLVM 的全部项目，也不要求从零实现所有 Toy 章节。

第 3 周：把图优化与量化约定接起来

本周问题：量化参数如何进入图，优化怎样保持轴、类型和误差约定？

本周材料：MLC 第 9 课图优化；MIT 第 5 讲：量化 I、第 6 讲：量化 II；MLIR 量化表示。论文从第十四节 AWQ 或 SmoothQuant 中选一篇，先读与本周实验对应的方法。回看本文第六节。

本周验收：换一个布局或分组方案时，知道哪些量化元数据也要变。设备尚不可用时完成图与数值验证；没有受支持的低比特内核时，性能项明确留待上板验证，不能用模拟量化的速度代表 NPU 低比特性能。

第 4 周：让 prefill、decode 与 KV 状态保持一致

本周问题：不同入口与形状变体怎样更新同一份逻辑会话状态？

本周材料：Core ML 状态模型指南、WWDC24 状态模型片段；按所选路径参考 ExecuTorch LLM 导出或 LLM-TPU 示例。示例用于理解接口，实验仍沿用本周可执行的后端。回看本文第五、十节。

本周验收：至少能复现并修正一个刻意引入的状态错误，例如 KV 写入偏移或掩码错误。后端不支持某种状态机制时，用显式缓冲区验证语义，并记录真实设备接口还缺什么；不要为完成全部案例而更换多套引擎。

第 5 周：从张量程序走到布局与缓冲区

本周问题：同一个数学计算，为什么不同分块、布局与存储安排会产生不同成本？

本周材料：MLC 第 2、3 课，按需回看第 9 课；MLIR Bufferization；CMU Data Layouts讲义；TPU-MLIR 官方索引中的 LayerGroup 教学。回看本文第七节。

本周验收：能够解释“局部更快而子图更慢”的一种可能原因，并用受控实验验证。MLC notebook 作为独立小练习即可，不要求同时深改 TVM 和 MLIR 两套编译器；CPU 分块实验建立机制理解，实际 NPU tile 与 DMA 结论仍由对应后端和设备验证。

第 6 周：读通一条目标后端，确认真实执行边界

本周问题：一个模型操作在哪一层被接受、改写、拒绝或交给其他设备？

本周材料：从第 13.5 节 A/B/C 中选一条；沿用第一周选定的设备与 SDK。厂商指南用于完成运行，源码与诊断信息用于定位决策。

本周验收：能根据证据说明哪些阶段在什么设备执行。没有真实 NPU 时，以可运行后端完成接入和编译机制实验，另附 NPU 适配差距清单；此时产出是可复现的编译学习项目，上板性能验收仍未完成。

第 7 周：围绕一个瓶颈做优化，而后决定是否保留

本周问题：哪一项编译决策值得改变，怎样证明收益来自这项变化？

本周材料：围绕已经观察到的瓶颈查对应 pass、内存、分区或厂商文档。自动调优有实际需求时再选看 MLC 第 5 课；新课程只用于解决当前问题。回看本文第十一、十二节。

本周验收：结果不要求达到某个加速倍数。能用证据推翻最初假设、解释编译器为何已经做过某项优化，或者发现一个明确的退化条件，都是合格的工程结论。

第 8 周：补稳定性，完成可复现交付

本周问题：另一位开发者能否在声明的条件下重建、运行并理解你的结果？

本周材料：所选工具链的产物兼容、运行时与测试说明；回看本文第十、十二节。停止新增大方向，把时间留给复现和解释。

本周验收：读者可以按说明重现关键行为，并知道哪些结论来自真实 NPU、哪些来自 CPU 或数值模拟。最终报告应同时保留失败尝试、未解决问题和下一步最值得投入的技术点。

13.5 第六周的三个深入方向，只选一条

A：模型接入与图编译。 优先从 ExecuTorch Qualcomm LLM 示例、ORT QNN EP或 LiteRT Android中选一个。先运行已有模型，再研究一个分解、量化参数或能力判断；对照修改前后的分区和执行位置。RKLLM 用户可以采用同样的接口分析方法，但厂商工具未开放的内部 pass 不能作为可修改范围。

B：编译器与硬件映射。 希望阅读较完整的 NPU 编译流程，可以选择 TPU-MLIR，沿导入、Top/Tpu IR、局部内存与代码生成中的一小段深入；使用 Intel 平台可先读 NPU Compiler 的 MLIR 入门和项目结构。具体改动限于一个可执行、可测试的规则，不以读完仓库为目标。

C：张量程序与运行时。 有受支持 AMD 设备和底层编程兴趣时，可以从 MLIR-AIE / IRON官方小示例研究计算、数据流和缓冲区；选择集成方向时，则在既定运行时中追踪状态创建、注册、执行和释放。前者深入硬件映射，后者深入执行接口，二者选一个具体问题即可。完整 LLM 的效果需要接回模型流程验证，小 GEMM 成功只能证明这条计算路径。

13.6 AI 怎样参与，才能真正缩短学习周期

适合交给 AI 的工作包括：解释已提供的 IR、导航固定版本源码、生成重复性测试输入、整理实验配置、比较两份 pass 输出，以及根据现有日志提出排查顺序。把当前版本、实际代码、完整错误和硬件条件一起提供，通常比泛泛要求“写一个 NPU 编译器”更有效。

三个可复用的提问模板：

 1
 2
 3
 4
 5
 6
 7
 8
 9
10

阅读这份变换前后 IR：逐项说明形状、类型、量化参数和状态变化。
列出该重写保持语义的前提，并给出至少一个不应匹配的反例。
对材料中不能确认的内容，明确指出需要继续查看哪处定义。

这是固定提交的 pass 与测试入口：请追踪匹配、合法性检查和替换。
提出最小改动及对应测试；不要使用该版本不存在的 API。

这是基线与改动的原始数据、设备和计时方法：
检查比较条件是否一致，区分事实、推测和仍缺少的测量，
并提出一次能够区分主要原因的后续实验。

每周最后一天，关掉 AI 辅助，用自己的话解释一次 IR 变化、一个失败输入和一项取舍。AI 生成的规则必须经过正反例与数值检查；没有运行过的命令、没有采集过的设备数据，不写成实验结果。

13.7 进度落后时，保护核心顺序

如果某周卡住，先缩小模型、图和形状集合，再减少选修内容；不要省略正确性、状态与真实执行位置检查来维持表面进度。学习路线的价值来自知识之间能互相解释，以及每一步都留下可检验的结果。

十四、进一步阅读：按工程问题选择论文和源码

14.1 核心材料与选读顺序

前面的课程负责搭起概念，下面的材料用于回答已经遇到的问题。每次阅读写下四件事：它解决什么约束、采用什么表示、依赖什么假设、怎样验证有效。

论文中的速度与模型范围依赖其设备、软件和评测条件。读过一种方法以后，先在自己的负载上找出对应约束，再决定是否复现；不要把发表年份或单项加速数字当作选型排序。

14.2 后续跟踪什么，暂时放下什么

建议持续跟踪四类变化：目标设备的模型/精度支持矩阵、所选编译器的 IR 与 pass 变化、状态和动态形状能力、编译产物与运行时兼容性。这些信息直接影响模型升级、性能和维护成本；遇到版本变化时，重跑已经建立的小模块、状态边界和完整模型检查。

当基础链路稳定后，再按真实需要扩展：长上下文和 KV 压缩、多模态前端、稀疏计算或混合专家（Mixture of Experts，MoE）的动态执行、自动调优和更高层的调度表达。它们各有价值，但进入学习主线的依据应是目标硬件与当前瓶颈，而非同时追逐所有新名词。

归纳起来，这张图谱的核心是：保留模型语义，表达硬件约束，验证程序变换，让编译产物与有状态执行正确衔接。 深入一套能修改、能观察、能验证的工具链，再用共同的问题比较其他实现，就能逐步形成面向端侧 NPU 的编译工程判断力。