第14章 TorchInductor：从 ATen IR 到 Triton kernel

“Inductor takes a graph of ATen ops and produces blazing-fast Triton kernels by being aggressive about fusion and conservative about correctness.”

—— Jason Ansel, “TorchInductor design” PyTorch Conference 2023

14.1 Inductor 在编译栈的位置

回顾整个 torch.compile 链路：

graph LR
    User["@torch.compile<br/>用户代码"]
    User --> Dy[Dynamo<br/>第 12 章<br/>FX Graph]
    Dy --> Aot[AOTAutograd<br/>第 13 章<br/>fw + bw 子图]
    Aot --> Ind[Inductor<br/>本章]
    Ind --> Tri[Triton kernel<br/>GPU]
    Ind --> Cpp[C++ + OpenMP<br/>CPU]

    style Ind fill:#fef3c7,stroke:#f59e0b,stroke-width:2px

Inductor 拿到的是一张 纯函数式的 ATen FX Graph（AOTAutograd 已经 functionalize 过），它要把这张图变成实际可执行的代码。源码在 torch/_inductor/，其中四个最大的文件：

• compile_fx.py (3047 行) — 编译入口
• lowering.py (7935 行) — ATen → Inductor IR
• ir.py (9966 行) — Inductor IR 节点定义
• scheduler.py (7158 行) — fusion 调度
• codegen/triton.py (6536 行) — Triton 代码生成

总规模在 v2.11 实测约 213,000 行 Python（v2.4 时仅 ~120k，v2.x 期间近乎翻倍），是 PyTorch 编译器栈最复杂、增长最快的部分。

14.2 Lowering：ATen → Inductor IR

打开 torch/_inductor/lowering.py，会看到几千个 @register_lowering 装饰器：

@register_lowering(aten.add)
def add(x, y, alpha=1):
    if alpha != 1:
        y = ops.mul(y, alpha)
    return ops.add(x, y)

每个 ATen 算子被注册一个 lowering 函数，把它翻译成 Inductor IR ops。Inductor IR 比 ATen 更底层 —— 它用循环 + 索引表达式描述计算，而不是 high-level 算子。

核心 IR 类（ir.py）：

关键设计：每个 IR 节点都用 indexing_fn(index) → 标量计算表达式 描述。比如 relu(x) 是一个 Pointwise，它的 inner_fn 是 lambda idx: ops.maximum(loader(idx), 0)。这种”循环 body 用纯函数描述”是 Inductor 实现 fusion 的基础 —— 两个 Pointwise 的 inner_fn 可以直接拼起来。

14.3 Scheduling：fusion 决策的核心算法

Lowering 结果是一组 IR 节点，每个节点是一段循环。Scheduler 决定哪些节点合并成一个 kernel：

节点 A: out_a[i] = relu(x[i])         # Pointwise
节点 B: out_b[i] = out_a[i] * 2       # Pointwise

→ 合并成: out_b[i] = relu(x[i]) * 2   # 一个 kernel, 一次内存读, 一次写

fusion 是 GPU 性能的最大杠杆。每次访存要几百纳秒，融合两个 op 省一次内存往返，对 memory-bound 算子（如逐元素操作）能 2x 加速。

graph LR
    subgraph NoFuse[不 fuse: 4 次 HBM 读写]
        X1[load x] --> A1[relu]
        A1 --> S1[store out_a]
        S1 -.HBM.-> L2[load out_a]
        L2 --> M1[mul 2]
        M1 --> S2[store out_b]
    end

    subgraph Fuse[fuse: 2 次 HBM 读写]
        X2[load x] --> A2[relu]
        A2 --> M2[mul 2]
        M2 --> S3[store out_b]
    end

    style NoFuse fill:#fee2e2
    style Fuse fill:#dcfce7

scheduler.py 实现的 fusion 决策考虑：

• 依赖关系：B 依赖 A → 可以 fuse
• 循环可对齐：A 是 [N] Pointwise、B 是 [N] Pointwise → 循环维度一致 → 可 fuse
• 寄存器压力：fuse 后中间变量数量是否爆炸
• 内存读写量：fuse 后总 byte 是否减少
• 共享中间值：A 输出被 B 和 C 共用 → fuse 哪个更优

最终算法是 贪心 + 启发式：从拓扑序开始，每次尝试把当前节点与前驱融合，估算收益（节省的访存 byte），收益正就 fuse。

FusedSchedulerNode（line 1866）表示融合后的节点，ForeachKernelSchedulerNode（line 2216）是 multi-tensor 优化的 _foreach_* 类操作专门融合节点。

14.4 Triton Codegen：把 IR 转成 Triton DSL

Triton 是 OpenAI 开发的 GPU DSL，比 CUDA C++ 更高级 —— 用户写”逻辑上的 block-level 循环”，Triton 编译器自动决定 thread mapping、内存合并、shared memory 等。

Inductor 生成的 Triton 代码大致长这样（简化版）：

@triton.jit
def fused_relu_mul_kernel(in_ptr0, out_ptr0, xnumel, XBLOCK: tl.constexpr):
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel

    tmp0 = tl.load(in_ptr0 + xindex, xmask)
    tmp1 = tl.maximum(tmp0, 0)        # relu
    tmp2 = tmp1 * 2                    # mul
    tl.store(out_ptr0 + xindex, tmp2, xmask)

codegen/triton.py:TritonKernel（line 2503）是核心生成类。它遍历 fused IR 节点的 inner_fn，把每条 Inductor IR op 翻译成 Triton 操作：

整个 codegen 是 string templating：拼接 Python 字符串组装 Triton 源码，再交给 Triton 的 triton.jit 编译成 PTX → CUBIN。

14.4.1 为什么用 Triton 而不是手写 CUDA

写一个 element-wise CUDA kernel 要决定：

• block size（一般 128 / 256 / 512，但最优值依 dtype / shape 变化）
• thread coalescing（让相邻 thread 访问相邻地址）
• shared memory 用法
• vectorized load（float4 / __half2）

每个算子都人工调这些参数太累。Triton 把这些决策全部交给编译器 —— 用户只写”block-level 循环”，Triton 自动选 block size、自动 vectorize、自动 coalesce。

这让 Inductor codegen 大大简化：只生成 logical Triton 代码，性能优化交给 Triton 编译器。这也是为什么 PyTorch 团队选择 Triton 而不是 LLVM —— 抽象层次更适合”我从 high-level IR 自动生成”的需求。

14.4.2 配 autotune 找最优 block size

Triton 还提供 @triton.autotune —— 给一组候选 block size，运行时跑一遍选最快的。Inductor 自动给生成的 kernel 加 autotune：

@triton.autotune(configs=[
    triton.Config({'XBLOCK': 64}),
    triton.Config({'XBLOCK': 128}),
    triton.Config({'XBLOCK': 256}),
    triton.Config({'XBLOCK': 1024}),
], key=['xnumel'])
@triton.jit
def kernel(...):
    ...

第一次跑 kernel 时 autotune 跑所有候选、记录哪个最快、缓存结果。后续 input shape 类似时直接用缓存最优配置。这套机制让 Inductor 在没人手动调参的情况下也能拿到接近手写的性能。

14.5 CPU 后端：C++ + OpenMP

GPU 上生成 Triton，CPU 上 Inductor 生成 C++ + OpenMP 代码。codegen/cpp.py 生成形如：

#pragma omp parallel for
for (int i = 0; i < N; ++i) {
    auto tmp0 = in_ptr0[i];
    auto tmp1 = std::max(tmp0, 0.f);
    out_ptr0[i] = tmp1 * 2;
}

加 #include <ATen/native/cpu/Vectorized.h> 让编译器自动 SIMD 化。整段 C++ 通过 torch.utils.cpp_extension 编译成 .so，运行时用 ctypes 加载。

CPU codegen 性能不如 Triton 给 GPU 的优化激进 —— 因为 C++ 编译器（gcc/clang）的优化水平不如 Triton 在 GPU 上的精细。这是 PyTorch 在 CPU 训练上比 GPU 加速比小的根本原因。

14.6 一个完整 trace：a + b * c

@torch.compile
def f(a, b, c):
    return a + b * c

Inductor 收到 ATen graph：mul(b, c) → add(a, _)。流程：

flowchart LR
    FX["fx Graph<br/>mul b c → add a _"]
    FX --> LO["Lowering<br/>每个 op 转 Pointwise IR"]
    LO --> IR["IR Nodes<br/>node mul: lambda i  load b i * load c i<br/>node add: lambda i  load a i + intermediate"]
    IR --> SC["Scheduling<br/>检查 add 依赖 mul  循环维度对齐<br/>决定 fuse 成一个 kernel"]
    SC --> CG["Codegen<br/>生成 Triton kernel<br/>load b/c → mul → load a → add → store"]
    CG --> KER["fused_mul_add 一个 GPU kernel<br/>4 次 HBM 访问替代 5 次"]

    style FX fill:#fef3c7
    style SC fill:#dbeafe
    style CG fill:#dcfce7
    style KER fill:#fce7f3

各步细节：

1. Lowering：mul → Pointwise(inner_fn=lambda i: load(b,i) * load(c,i))，add → Pointwise(inner_fn=lambda i: load(a,i) + load(intermediate,i))
2. Scheduling：发现 add 依赖 mul，循环维度都是 [N]，可 fuse
3. Codegen：生成一个 Triton kernel:

@triton.jit
def fused_mul_add(a_ptr, b_ptr, c_ptr, out_ptr, xnumel, XBLOCK: tl.constexpr):
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)
    xmask = xindex < xnumel

    tmp0 = tl.load(b_ptr + xindex, xmask)
    tmp1 = tl.load(c_ptr + xindex, xmask)
    tmp2 = tmp0 * tmp1
    tmp3 = tl.load(a_ptr + xindex, xmask)
    tmp4 = tmp3 + tmp2
    tl.store(out_ptr + xindex, tmp4, xmask)

——三个原始 ATen 算子（mul / add）变成一个 GPU kernel，内存读写从 5 次减到 4 次，性能 GPU memory bandwidth bound 提升 ~25%。

14.6.5 fx_passes：Lowering 之前的图变换

torch/_inductor/fx_passes/（28506 行）有十几个 Pass 在 Lowering 之前对 FX Graph 做改写：

每个 Pass 是一个独立的图变换，按 compile_fx.py 里固定顺序跑。这些 Pass 让 Inductor 拿到的不是”AOTAutograd 原始输出”，而是已经做过几轮高级优化的图 —— 后面的 Lowering / Scheduling 在更优的 IR 上工作。

reinplace 特别有意思：第 13 章 §13.4 提过 AOTAutograd 把 inplace 全部 functionalize 掉，让中端图分析能干净做。但最终 codegen 出 Triton 时，inplace 才能省内存。reinplace Pass 在 Lowering 前把”逻辑上独立但实际可复用”的张量重新 inplace。这是”先 functionalize 简化分析、再 reinplace 恢复性能”的工程闭环。

14.6.6 Pattern Matcher：自动重写为 fused 算子

pattern_matcher.py（2376 行）实现了 Inductor 的”模式匹配重写”。它能识别 FX Graph 里特定子图，整段替换成更高效的算子调用。

经典例子：SDPA pattern matching。用户写：

# 朴素 attention 实现
scores = q @ k.transpose(-2, -1) / sqrt(d)
attn = F.softmax(scores, dim=-1)
out = attn @ v

Pattern matcher 识别这一段，自动重写成 aten._scaled_dot_product_flash_attention（FlashAttention-2 实现）。用户什么都不用做就拿到 FlashAttention 加速。

类似地有：

• LayerNorm pattern → fused aten.native_layer_norm
• Conv + BN + ReLU → fused conv-bn-relu kernel
• Embedding lookup + scale → fused kernel

每个 pattern 是一段”匹配子图模板 + 替换函数”。Pattern matcher 用 graph isomorphism 算法在大图里找子图。这套机制让 Inductor 不用每个用户手写 attention，自动用上 FlashAttention。

14.6.7 AsyncCompile：多进程编译加速

Triton 编译一个 kernel 需要几百毫秒到几秒（PTX → CUBIN）。一个大模型 forward + backward 可能产生几百个 kernel —— 串行编译要几十分钟。

async_compile.py（755 行）实现了多进程并行编译：

• 主进程生成 Triton 源码 + 计算 hash
• 派给一个 worker 进程池（几十个 worker）
• 每个 worker 调 Triton compiler 编译
• 主进程并发等待所有 future
• 所有 kernel 编译完才返回最终的 compiled callable

实测能让”几十个 kernel 编译”从 2 分钟降到 5-10 秒。这是 v2.x 之后 torch.compile 启动时间能从”几分钟”压到”几十秒”的根本。

worker 进程间通过 hash 共享缓存：同一个 Triton 源码（不同 input shape 但生成的代码一样）被认成同一份编译产物，避免重复编译。

14.6.8 Wrapper Codegen：粘合 Python 与 Triton

Inductor 生成的不只是 Triton kernel，还有包装这些 kernel 的 Python 代码（codegen/wrapper.py）。简化后产物形如：

# Inductor 生成的 wrapper 函数
def call(args):
    arg0_1, arg1_1 = args
    args.clear()
    buf0 = empty_strided_cuda((1024, 1024), (1024, 1), torch.float32)
    triton_per_fused_relu_mul_0.run(arg0_1, arg1_1, buf0, 1048576, grid=grid(1024))
    del arg0_1, arg1_1
    return (buf0,)

这段 Python 代码包了 Triton kernel launch 与中间 buffer 分配。它本身也被 Python 解释器跑，但因为内部都是 Triton kernel 的 batched launch，Python 端开销可忽略。

codegen/cpp_wrapper_cpu.py 与 cpp_wrapper_gpu.py 是更激进的版本：把 wrapper 也编译成 C++（避免 Python 解释器开销），让”零 Python overhead 的部署”成为可能。这是面向 mobile / edge 部署的关键。

14.6.9 GEMM Template：matmul 的专用 codegen

第 14.7 节会讲 ExternKernel（cuBLAS / cuDNN 直接调用），但 Inductor 还有第三条 GEMM 路径：自家 GEMM template。torch/_inductor/kernel/mm.py:305 tuned_mm 是入口：

def tuned_mm(mat1, mat2, out_dtype=None, *, layout=None):
    # 收集所有候选 GEMM 实现
    choices = []
    if use_aten_gemm_kernels():
        choices.append(aten_mm.bind(...))                # cuBLAS / cuDNN
    if use_triton_template(layout):
        for config in mm_configs(...):
            choices.append(mm_template.maybe_append_choice(config))   # Triton GEMM
    if use_cutlass_template(layout):
        choices.append(cutlass_template.maybe_append_choice(...))     # CUTLASS

    # autotune 在所有候选里选最快的
    return autotune_select_algorithm("mm", choices, ...)

——matmul 调用时 Inductor 同时生成 cuBLAS / Triton / CUTLASS 三种实现，autotune 实测选最快的。这与普通 Pointwise 算子（直接 codegen Triton）路径完全不同。

为什么 GEMM 要这样？因为 matmul 性能差异巨大：

• cuBLAS 在大 matmul（如 4096×4096）上接近峰值，但 launch overhead 高（~10us）
• Triton GEMM template 可以与 epilogue（如 add bias、relu）fuse，端到端更优；但小 matmul 上不如 cuBLAS
• CUTLASS 在某些特殊 dtype（fp8）上是唯一选择

mm.py 还有几个特殊 template（:216 DecomposeKSugraphTemplate / :254 ContiguousTemplate）：

• DecomposeKSugraphTemplate：把大 K 维 matmul 拆成多个小 K + accumulate（避免 K 维太大让单 kernel reduction 慢）
• ContiguousTemplate：只处理 contiguous input 的快路径

整套机制让 LLM 训练 / 推理里的 attention（QKV 投影 + softmax 后 mm）能 自动选最优 GEMM 实现 + fuse epilogue。max_autotune_gemm=True（§15.6.6）让这套机制穷举所有 config，能再多 5-10% 性能。

cpp_gemm_template.py 是 CPU 端 GEMM template，1828 行实现 oneDNN / mkldnn 集成 + 自家 fp32 / bf16 GEMM 微 kernel 的 C++ codegen。CPU 上 LLM 推理的核心性能优化集中在这里。

14.6.10 IR 节点的完整层次：Buffer / ComputedBuffer / TensorBox / StorageBox

§14.2 提到 Inductor IR 的核心是 Loops（Pointwise / Reduction），但实际 IR 是个有层次的对象图。打开 ir.py（9966 行），会看到几个不同抽象层级的类：

调用关系：TensorBox 包 StorageBox，StorageBox 包 Buffer，ComputedBuffer 持有 Loops（具体计算）。这种”四层包装”是 PyTorch 内部 Tensor 的镜像（第 2 章 §2.2 的 Tensor / TensorBase / TensorImpl / Storage 三件套思想）。

graph TB
    TB[TensorBox<br/>用户视角的 tensor]
    SB[StorageBox<br/>共享存储]
    Buf[Buffer<br/>具体内存 buffer]
    CB[ComputedBuffer<br/>需要计算才能填的 buffer]
    Loops[Pointwise / Reduction / Scan<br/>计算定义]
    Layout[FixedLayout / FlexibleLayout<br/>strides + offset]

    TB --> SB
    SB --> Buf
    Buf -.子类.-> CB
    CB --> Loops
    Buf --> Layout

    style TB fill:#fef3c7
    style SB fill:#dbeafe
    style Buf fill:#dcfce7
    style Loops fill:#fce7f3

为什么要这么多层？因为 Inductor 要在不同阶段做不同变换：

• StorageBox 让 view 操作（reshape / view）零成本 —— 多个 TensorBox 共享同一个 StorageBox，类似 ATen 的 view
• FlexibleLayout 表示”layout 还没决定，编译器可以选最优 strides”，到后期 decide_layout() 才固化
• MultiOutputLayout 用于一个 kernel 多输出（如 batch norm 同时输出 mean / var / output）

理解这套层级，再读 ir.py 就不会被几千行类定义绕晕 —— 每层的职责清晰、互不重叠。

14.6.11 Memory Planning：buffer 复用与生命周期

torch/_inductor/memory.py（1108 行）实现了 Inductor 的”内存复用”分析。基本原理：

Buffer A: alive [step 1, step 5]
Buffer B: alive [step 6, step 9]  ← B 完全在 A 死亡之后

→ B 可以复用 A 的内存地址 (assign 同一段)

memory_planning 算法（Lin’s algorithm 改编）做的事：

1. 对每个 buffer 算”出生时间”（首次写入）和”死亡时间”（最后读取）
2. 按死亡时间排序
3. 后续 buffer 优先复用已死的 buffer 内存
4. 最终所有 buffer 的 peak memory 显著低于”每个 buffer 独立分配”的 baseline

对一个 transformer 训练 step，这套优化能让 activation 显存峰值降 30-50%。Inductor 默认开启，用户感受不到但底层在自动节省。

scheduler.py:3011/3453 调 assign_memory_planning_info_for_scheduler_buffers 把每个 SchedulerBuffer 标注 reuse 关系。最终 codegen 出的 wrapper 代码里会看到：

buf3 = empty_strided_cuda(...)
buf3 = reuse_storage_of(buf0)    # buf3 复用 buf0 已死的内存

这种”显式 reuse”让 Triton kernel 不需要新 alloc，直接写入复用 buffer。第 4 章 §4.x 的 caching allocator 在 Inductor 路径上几乎不起作用 —— 因为 Inductor 在编译期就做完了内存规划。

14.6.12 FxGraphCache：跨进程编译产物缓存

codecache.py:1206 的 FxGraphCache 让 Inductor 编译产物持久化到磁盘：

# 第一次跑: 编译 + 写盘
torch.compile(model)(x)    # Inductor 编译, 写到 ~/.cache/torch/inductor/...

# 重启进程后跑同样代码
torch.compile(model)(x)    # cache hit, 跳过编译, 直接 load .so

cache key 怎么算？FxGraphCachePickler (:485) 把 FX Graph + example_inputs shape/dtype + Inductor config + PyTorch version 一起 hash。任何一项变化就 cache miss，避免拿到错误的旧产物。

对生产推理服务，FxGraphCache 让”重启服务首次请求”从 30 秒（重新编译）降到 < 1 秒（load .so）。设 TORCHINDUCTOR_CACHE_DIR=/path/to/persistent 让 cache 跨容器共享 —— 同一个模型多个 pod 可以共用编译产物。

cache 失效的常见原因：

• 改了模型架构（FX Graph 变化）
• 升级了 PyTorch 版本
• 改了 _inductor.config 任何 flag
• 输入 shape 偏离了 trace 时的假设

第一次部署时预编译 + warmup 能避免线上首请求慢，是大模型推理服务的标准操作。

14.6.13 ChoiceCaller / autotune select_algorithm

torch/_inductor/select_algorithm.py 的 ChoiceCaller 是 GEMM template autotune（§14.6.9）的核心。每个 ChoiceCaller 是”一种算法实现 + 它的 benchmark 接口”：

class ChoiceCaller:
    def benchmark(self, *inputs):
        # 跑这个候选算法 N 次, 返回平均时间
        ...

    def call(self, *inputs):
        # 实际调用此算法, 返回结果
        ...

autotune_select_algorithm("mm", choices, ...) 内部：

1. 对每个 choice 跑几次 benchmark（用真实 input shape，不是假数据）
2. 选最快那个
3. 把”shape -> 最优 choice”缓存到 disk
4. 后续相同 shape 的 mm 直接命中缓存

ExternKernelCaller（:2438）是接 cuBLAS / cuDNN 等外部 kernel 的 ChoiceCaller 实现。用户也能注册自己的 ChoiceCaller —— 比如某厂商把自家硬件的 GEMM kernel 包成 ChoiceCaller，让 Inductor autotune 能选到它。

这套机制让 Inductor 在 GEMM 等关键算子上性能自适应 —— 用户不调参，编译器自己选最优。

14.6.14 Subgraph 编译与 HigherOrderOperator

某些用户代码包含”子图”结构 —— 如 torch.utils.checkpoint（第 7 章 §7.5.3）、while_loop / cond 控制流、flex_attention 自定义 attention。Inductor 把这些子图当作独立编译单元：

# checkpoint 内部的 forward 是一段 subgraph
def forward(x):
    out = checkpoint(my_block, x)   # my_block 是个子函数
    return out

# Inductor 把 my_block 编成独立的 subgraph compiled callable
# 主图调用时通过 HigherOrderOperator 接进去

HigherOrderOperator (HOP) 是 PyTorch v2.x 引入的概念：算子的”参数”可以是另一段计算图。torch.cond、torch.while_loop、torch._higher_order_ops.flex_attention 等都是 HOP。Inductor 给每个 HOP 单独 lower 它的子图。

这套机制让 可微分 attention（如 FlexAttention） 等高级用法在 Inductor 路径上能完整编译，而不是 fall back 到 eager。第 22 章 §22.6.5 的 opcheck test_aot_dispatch_* 也会验证 HOP 路径。

14.6.15 Scheduler 的 fusion 算法细节

§14.3 给了 Scheduler 的高层视角，这里展开 scheduler.py（7158 行）的具体 fusion 算法。核心两个函数：

can_fuse(node1, node2) -> bool：判断两个节点能否融合。决策树：

1. 数据依赖：node2 依赖 node1 输出，且没有循环依赖
2. 循环维度可对齐：两节点的迭代空间能映射到同一组 loops（最常见的”同 shape”特例）
3. dtype 兼容：fuse 后中间值的 dtype 不会让寄存器溢出
4. reduction 兼容：两个 reduction 维度一致才能 fuse 成一个 kernel
5. 无 atomic 阻塞：写到同一 buffer 的多个写入必须都通过 atomic

任意一条不满足就拒绝 fuse。第 4 / 5 条是 reduction 类算子（如 sum / mean）经常无法 fuse 的根本原因。

score_fusion(node1, node2) -> int：给可 fuse 对打分。分数构成：

• 节省的内存读写：基础项。fuse 后省掉的 byte 数（中间 buffer 不写也不读）
• reduction 边界 bonus：reduction → pointwise fuse 收益大（reduction 输出小，pointwise 跟在后面几乎免费）
• broadcast penalty：广播会让 fuse kernel 的循环展开复杂、寄存器压力大
• co-located 同 device：必须同 device 才有意义

Scheduler 用 贪心 + 优先队列 应用 fusion：每次取出分数最高的可 fuse 对，融合，重新评估剩余对。这种局部贪心不保证全局最优但工程上够用 —— 对大多数 transformer 训练能拿到 80-90% 的最优 fusion 收益。

scheduler.py:1866 的 FusedSchedulerNode 是融合后的产物。一个 FusedSchedulerNode 可能包几个原始 SchedulerNode 的 inner_fn，运行时它们一起进同一个 Triton kernel。

14.6.15.1 Foreach Fusion

scheduler.py:2216 的 ForeachKernelSchedulerNode 是特殊融合：把多个独立但 shape 相似的算子（如 _foreach_add_(params, grads) 处理几百个张量）一次 launch。第 10 章 §10.6.1 提过 torch._foreach_* 算子，Inductor 在编译路径上对它们做单独优化 —— 多张量并行处理在 Inductor codegen 出来的是单个 kernel + grid 维度按张量数展开。

14.6.16 CSE：消除冗余索引计算

Inductor 的 IR 大量使用索引表达式（如 idx0 = pid * BLOCK + offset、mask = idx0 < n）。同一个表达式可能在 kernel 不同位置重复出现。CSE（Common Subexpression Elimination）让相同表达式只算一次。

具体在 codegen/common.py 的 CSE 类。每次生成代码时：

# 朴素 codegen 可能产出
tmp0 = tl.load(x + idx)      # idx = pid*BLOCK + arange
tmp1 = tl.load(y + idx)      # 又算一遍 idx

# CSE 后
idx_cse = pid*BLOCK + arange
tmp0 = tl.load(x + idx_cse)
tmp1 = tl.load(y + idx_cse)

实测 CSE 能让 Triton kernel 寄存器压力降 20-30%，对寄存器受限的复杂 kernel（如大 K 的 GEMM）显著提升性能。

IndexExprCSE 是 CSE 的索引表达式特化版 —— 它知道索引表达式的代数恒等式（(i+0) == i、(i*1) == i、模运算化简等），能做比通用 CSE 更激进的化简。这套优化是”Triton 编译器拿不到、必须 Inductor 做”的工程价值 —— Triton 看到的是 IR 里已经化简过的表达式。

14.6.17 CPU 后端的 Vectorized 抽象

§14.5 提过 CPU 用 C++ + OpenMP，这里展开 SIMD 细节。codegen/cpp.py（CPU codegen 主文件）大量用 PyTorch 自家的 Vectorized<T> 抽象（aten/src/ATen/cpu/vec/）：

// Inductor 生成的 (简化) CPU kernel
#pragma omp parallel for
for (int64_t i = 0; i < N; i += Vectorized<float>::size()) {
    auto a = Vectorized<float>::loadu(in_a + i);
    auto b = Vectorized<float>::loadu(in_b + i);
    auto c = a * b + a;       // 自动向量化
    c.store(out + i);
}

Vectorized<float> 在 x86 是 AVX2/AVX512 寄存器（8/16 元素）、ARM 是 NEON（4 元素）、Apple Silicon 是 Accelerate framework。同一份 codegen 代码在不同平台编译出最优 SIMD，无需 Inductor 写多套。

特别值得提的是 bf16 / fp16 在 x86 上的处理：x86 没有原生 fp16 指令（除了最新 Sapphire Rapids），Vectorized 通过”fp16 → fp32 → 计算 → fp16”的 cast 链实现。Inductor codegen 出来的 cast 会被 C++ 编译器进一步优化成最少的 cast 指令。

CPU codegen 还有专门的 cpp_gemm_template.py（§14.6.9）：CPU GEMM 用 oneDNN 或者自家 micro-kernel（cpp_micro_gemm.py）。LLM 在 CPU 推理时这条路径承担核心性能。

14.6.18 inductor.config 的 200+ flags 概览

torch/_inductor/config.py 700+ 行注册了几百个 config flag。按功能分组：

每个 flag 都对应一个具体的 codegen / 优化 pass。生产级调优要熟悉这 200+ flag 才能在每个场景找最优组合。但 90% 用户用默认值即可 —— PyTorch 团队把”通用最优”调到 default。

config.fallback_random 是一个有意思的 flag：默认 True，意思是”遇到 random 算子（torch.rand 等）让它走 ATen，不走 Inductor 的 functional random”。这是因为 Inductor 的 random 实现还不完善；用户配 False 强制走 Inductor，能享受 fuse 但要承担”重 forward 时随机数不一致”的风险（第 13 章 §13.4 提过的 random functionalize）。

14.6.19 一段 LayerNorm 的完整 Inductor 编译产物

把整章串起来，看 nn.LayerNorm 在 Inductor 下编译成什么。源码：

@torch.compile
def f(x, weight, bias):
    return F.layer_norm(x, [768], weight, bias, 1e-5)

经过完整链路：

1. Dynamo trace：捕获 fx graph，含 aten::native_layer_norm 一个节点
2. AOTAutograd functionalize：layer_norm 没 inplace，跳过
3. AOTAutograd decomposition（§13.6.8）：把 native_layer_norm 拆成 mean → var → rsqrt → sub → mul → add 共 8 个基础节点
4. Inductor Lowering：每个基础节点变成 IR 节点（Pointwise / Reduction）
5. Scheduler fusion：mean + var 是 reduction，但能 fuse 在同一 reduction kernel；后面 sub/mul/add 串成 pointwise，作为 reduction 的 epilogue 一起 fuse
6. Triton codegen：生成单个 Triton kernel，含 reduction 部分（mean + var 在 shared memory）+ epilogue（normalize 部分）
7. Triton 编译：Triton DSL → MLIR → PTX → CUBIN
8. 运行时：Wrapper 调用此 kernel + autotune 选最优 BLOCK 配置

最终一行 F.layer_norm(...) 调用变成一次 Triton kernel launch，性能与 PyTorch 内置 fused layer_norm（cuBLAS / cuDNN 提供）几乎相当，有时更快（因为能继续与周围 op fuse —— 比如 layer_norm + linear 在某些场景能融成单 kernel，cuBLAS 给的 fused layer_norm 做不到这个）。

理解这条全旅程，你就理解了 torch.compile 在 Transformer 训练上的真实加速来源 —— 不是某个魔法算法，是端到端编译链的累积优化。

14.7 Inductor 的”非 Pointwise”路径

复杂算子（matmul、conv、SDPA）不容易自动 fuse 出最优代码，Inductor 不自己生成，而是调用 现成的 cuBLAS / cuDNN / FlashAttention kernel。这些走 ExternKernel（IR 类型之一）路径，Inductor 只决定 layout 与 stride 转换，不生成 GEMM 代码。

但 fused attention 等场景，Inductor 在 aten._scaled_dot_product_flash_attention 调用前后能融合 layernorm / dropout，整体优于 eager。这种”自动选最优 GEMM 实现 + 周边 fuse”是 Inductor 在 transformer 训练上能 1.5-2x 加速的来源。

14.7.5 编译错误与 fallback 机制

Inductor 不是万能的。某些算子组合编译会失败 —— 复杂 indirect indexing、动态 shape 边界、外部库依赖等。torch._inductor 有完善的 fallback 体系：

flowchart TB
    Try[尝试 Inductor 编译]
    Try --> OK{成功?}
    OK -->|是| Triton[生成 Triton kernel]
    OK -->|否| Cause{失败原因?}

    Cause -->|某 op 无 lowering| Per[per-op fallback<br/>这个 op 用 ATen, 其他用 Triton]
    Cause -->|图整体 trace 失败| Graph[整图 fallback 到 eager]
    Cause -->|内存不够编译| Smaller[调小 BLOCK, 重试]

    style Per fill:#fef3c7
    style Graph fill:#fee2e2

具体机制：

• @register_lowering(aten.xxx, type_promotion_kind=...) 缺失 → Inductor 报 MissingOperatorWithoutDecomp，自动 fallback：把这个 op 留给 ATen 跑，前后的 Triton kernel 把数据传过来
• shape 推导失败 → 整图 fallback：TORCH_LOGS=dynamo 会看到 “graph break (inductor failed)”
• kernel 编译时间过长 → 触发 timeout，降级到不那么激进的 fusion 配置

这套 fallback 让 torch.compile 在生产代码里几乎不会让程序崩 —— 最坏情况 fallback 到 eager，性能不如 compile 但功能正常。这种”优雅退化”是 PyTorch v2.x 设计哲学的体现：编译是优化、不是必需。

14.7.6 自定义算子的 Inductor lowering 注册

§22.6 讲过 torch.library.custom_op 写自定义算子。如果想让自定义算子也走 Inductor 编译（不只是当不透明 op），需要注册 lowering：

from torch._inductor.lowering import register_lowering, ops

@register_lowering(my_lib.my_mul.default)
def my_mul_lowering(x, y):
    # 用 Inductor IR 重新表达 my_mul 的语义
    return ops.mul(x, y)

这种”用 Inductor IR 重新表达”让自定义算子能与周围算子 fuse，而不是当不透明 op 阻断 fusion。代价是要写两套实现（一份 eager kernel + 一份 Inductor lowering），但对性能敏感的核心算子值得。

实际国内 AI 芯片厂商接 Inductor 时，通常给自家硬件实现一套 codegen backend，复用 Inductor 的 lowering 框架。这是除了第 5 章 dispatcher 之外的另一条接入路径，让国产芯片能享受 Inductor 的 fusion / autotune 优化。

14.7.7 Triton kernel 的完整生命周期

讲了 Inductor 这么多，最后看一下产物（Triton kernel）从生成到运行的完整旅程：

1. Inductor 生成 Triton DSL (Python 代码)
   ↓
2. Triton compiler: Triton DSL → Triton IR (MLIR-based)
   ↓
3. Triton lowering: Triton IR → MLIR LLVM Dialect
   ↓
4. LLVM: → PTX (NVIDIA) / AMDGPU IR (AMD)
   ↓
5. NVCC ptxas / AMD lld: → CUBIN (NVIDIA) / HSA code object (AMD)
   ↓
6. CUDA driver / HIP runtime: load → 运行时 launch

整套链路在 torch/_inductor/async_compile.py（§14.6.7，class AsyncCompile 在 line 232）的 worker 进程里跑。每一步都是开源工具：

• Triton compiler 是 OpenAI 开源（独立项目 github.com/triton-lang/triton）
• MLIR 是 LLVM 子项目
• PTX/CUBIN 是 NVIDIA 闭源但有公开规范

CUBIN 二进制最终被 CUDA driver 加载到 GPU。torch._inductor 的 wrapper 代码持有 cudaModuleLoadData 返回的 module handle，通过 cudaLaunchKernel 触发执行。这套链路完全跳过 PyTorch 的 dispatcher（第 5 章）—— Inductor 编译产物是裸 CUDA kernel，运行时零 dispatcher 开销。

理解这条链路你就明白：torch.compile 的”加速”本质是把”用 Python 解释 + dispatcher 调度”的运行时开销，提前到编译期消化掉。运行时只剩 GPU 计算本身。

14.7.8 Inductor 与 Pattern Matcher 在 attention 路径上的协奏

LLM 训练 / 推理最热的是 attention。Inductor 在 attention 上做的事最丰富，把多个机制综合起来：

1. Pattern matcher 识别：用户写的”naive attention”代码（QK / sqrt / softmax / V）被识别成 SDPA 模式
2. 替换为 aten._scaled_dot_product_flash_attention：调 FlashAttention-2 实现
3. Epilogue fusion：attention 后的 dropout / linear projection 被 fuse 进 attention kernel 末尾
4. GEMM template 选择：QKV 投影矩阵的 qkv = x @ Wqkv 用 GEMM template autotune 选 cuBLAS / Triton / CUTLASS 最优
5. Memory planning：QKV 中间张量、attention scores 的内存复用

整套综合优化让 Inductor 在 transformer 训练上能拿到 1.5-2x 加速（相对 eager），而且几乎不需要用户手写优化 —— 写最朴素的 attention 代码就能享受 FlashAttention + fused epilogue + 最优 GEMM。

这就是 Inductor 真正的工程价值：让用户专注算法、编译器负责性能。这与传统 CUDA C++ “性能必须靠手写优化” 的模式形成鲜明对比，是 PyTorch 团队为 LLM 时代铺的工程基础。

14.8 几条工程经验

实战 Inductor：

1. TORCH_LOGS=inductor,output_code 看生成的 Triton 源码：能看到 fusion 决策、kernel 数量、autotune 结果

2. fusion 失败常见原因：循环维度不一致（如 [N] 和 [N, M]）、有 reduction 隔断、有 stride 不匹配的 view

3. 第一次跑会编译几秒到几分钟：Triton 编译 + autotune。设 TORCHINDUCTOR_CACHE_DIR 落盘，下次进程直接用

4. dynamic shape 多时性能下降：Inductor 对每个 shape 生成不同 kernel。设 dynamic=True 让一份 kernel 处理多 shape

5. 自定义 op 要注册 register_lowering：否则 Inductor 看到不认识的 op 触发 graph break，回退 eager

6. CPU 训练加速比小：Inductor 在 CPU 上生成 C++ + OpenMP，相对 ATen 直接调底层库，加速通常只 10-30%

7. mode='reduce-overhead' 启用 CUDA Graph：把多个 kernel launch 合成一次 graph replay，对小 batch 推理巨大加速

14.8.5 历史决策：为什么不直接生成 CUDA C++

Inductor 早期（v1.13 prototype）试过两条 codegen 路径：CUDA C++ 与 Triton。最终选 Triton 是几个工程考虑：

CUDA C++ 路径的问题：

• 每个 kernel 要手写 launch config（block / grid / shared memory size）
• 跨 GPU 架构（A100 / H100 / L4）调优要分别写
• 编译路径长（NVCC 慢、cubin 不可移植）
• 与 PyTorch C++ 端集成要 build .so，部署复杂

Triton 路径的优势：

• block-level 抽象让 codegen 不用管 thread mapping
• 跨架构统一编译（Triton 自动选最优）
• Python 端直接 jit，无 build 步骤
• autotune 框架成熟

代价是 Triton 不支持某些”激进 GPU 编程模式”（如 warp specialization、tensor memory accelerator on H100）。Inductor 在这些场景下走 ExternKernel（如 FlashAttention 用专门的 CUDA kernel 不走 Triton codegen）。这是”通用 codegen + 关键 op 走专用实现”的混合策略。

理解这条决策，你就明白为什么 Inductor 不试图与 cuBLAS 在 GEMM 上完全替代 —— 通用 codegen 路径在 GEMM 这种”几十年优化的算子”上很难超越专用库。Inductor 的价值在 fusion + 通用算子的 codegen，不在 reinvent GEMM。

14.8.6 Halide 思想的影响

Inductor IR 的”compute / schedule 分离”思想直接借鉴 Halide（image processing DSL，2012 年 MIT）。Halide 的核心范式：

• Compute：定义”算什么”（数学公式）
• Schedule：定义”怎么算”（loop ordering / vectorize / parallelize）

Inductor 对应到：

• Loops.inner_fn：compute 部分 —— 描述循环 body 怎么算
• Scheduler 的 fusion 决策：schedule 部分 —— 决定哪些 loop 合并、什么顺序

这种分离让”数学语义”与”性能调优”解耦。Halide 用户可以写一份 compute、试多种 schedule 找最优。Inductor 把”试 schedule”自动化（autotune），用户连 schedule 都不用写。

这条思想也影响了 TVM、MLIR / IREE、Triton 等其他编译器。Inductor 不是凭空发明，而是站在 Halide / TVM / Triton 几十年研究的肩膀上，把这些思想集成到 PyTorch 生态。

14.8.7 dynamic shape 在 Inductor 的处理

§12.5 提过 Dynamo 的 guards 处理 dynamic shape。Inductor 这一层也有相应支持：

# 用户指定 batch dim 是动态的
torch._dynamo.mark_dynamic(input, 0)
output = compiled_model(input)

Inductor 收到的 graph 里 batch dim 是 SymInt（符号整数），不是具体值。codegen 出来的 Triton kernel 把 batch dim 作为 kernel 参数，运行时传入：

@triton.jit
def fused_kernel(x_ptr, out_ptr, B: tl.constexpr, M: tl.constexpr, ...):
    # B 是 batch dim, 运行时传入
    ...

tl.constexpr 让 Triton 在第一次见到 B 的具体值时编译特化版。第二次相同 B 命中编译缓存；不同 B 重新编。这是 dynamic 与 specialized 的折中：B 维确实变化时不需要重 trace，但每个独特 B 值还是各编一次 kernel。

这套机制让 LLM 推理（不同 batch size、不同 sequence length）的 dynamic shape 路径在 Inductor 下能正确编译，而不是退回 eager。代价是 Triton 编译缓存项数会增加 —— 长跑服务里要监控缓存大小，必要时设 cache_size_limit 防膨胀。

14.8.8 Inductor 之外的两条编译路径

虽然 Inductor 是 torch.compile 默认 backend，PyTorch 主仓还有两条编译实验路径：

• aot_eager：只跑 AOTAutograd 不上 Inductor，主要用于调试。Inductor 编译出问题时切到这个 backend 看是不是 AOT 阶段错的
• TorchScript (torch.jit.script / torch.jit.trace)：v1.x 时代的编译路径，已被 torch.compile 取代但仍在维护。某些 mobile 部署场景（不能跑 Triton）仍用 TorchScript

新代码不要用 TorchScript —— 它的 graph trace 不稳定、不支持 dynamic shape、与 Inductor 路径不兼容。torch.export + AOTI（§15.6.7）是 v2.x 的统一部署方向。

14.9 跨书关联

• 《Rust 编译器之路》编译器后端 codegen：Inductor 的 IR-Lowering-Codegen 三段式与 LLVM 的 IR-MachineIR-Asm 完全同构
• 《vLLM 内核探秘》第 8 章 model runner：vLLM 也用 torch.compile + Inductor 编译模型 forward。理解本章能帮你解释 vLLM 启动慢的原因（首次 Inductor 编译）
• 《Triton 论文》（Tillet et al., 2019）：理解 Inductor 必读 Triton 原理。Triton 把”GPU 编程”抽象到 block 级别让 Inductor 这种自动 codegen 成为可能

14.9.5 Inductor 性能数字：实测加速比

具体数字（H100，PyTorch v2.11，典型 forward + backward）：

decode 加速比最大是因为 dispatcher 开销占比高（小算子密集）；大模型训练加速比相对小是因为 compute-bound，已经接近硬件峰值。

Inductor 不能 100% 替代手写优化：FlashAttention v3、cuDNN convolution、cuBLAS GEMM 等”上百人月优化”的 kernel 仍然手写更优。Inductor 的价值在 算子之间的胶水（fusion）和 没人专门优化的小算子。这两块 Inductor 做得足够好，让大模型训练 / 推理整体性能逼近 hand-optimized 水平。

14.9.5.5 Inductor 与 profiler 的协作

第 21 章讲过 PyTorch profiler。Inductor 编译产物在 profile chrome trace 里有特殊命名规则：

triton_per_fused_relu_mul_0    ← Inductor 生成的 fused kernel
                              fused_<op1>_<op2>_<id> 描述被 fuse 的算子

这个命名让你看 trace 时能立刻识别”这个 kernel fuse 了哪些 op”。如果你看到 triton_per_fused_layer_norm_dropout_linear_3 这种长名字，说明三个算子被 fuse 成一个 kernel，性能优秀。如果看到很多独立的 triton_per_fused_relu_0 / triton_per_fused_mul_1，说明 fusion 没起作用 —— 该排查为什么 scheduler 拒绝了这些 fuse（往往是循环维度不对齐或 reduction 隔断）。

TORCH_LOGS=output_code 让 Inductor 把生成的 Triton 源码 + wrapper Python 都打印到日志，结合 profiler 的 kernel 时间，能精确定位”哪段编译产物慢”。这套调试链路是 Inductor 性能调优的标配，比单纯看 chrome trace 信息丰富得多。

14.9.5.6 Inductor 与 torch.export 协作

torch.export(model) 把模型 trace 成 ExportedProgram（静态图，更严格）。AOTI（§15.6.7）就是基于 ExportedProgram 调 Inductor 编译。但有些算子在 torch.export 路径下需要特殊处理：

• HigherOrderOperator (§14.6.14)：control flow / checkpoint / flex_attention 等
• Custom op 必须有 register_fake（§22.6.5）：否则 export 时 shape 推导失败
• Effect token op：print / collective ops 在 ExportedProgram 里要保留 side effect ordering

torch._inductor 暴露的 aoti_compile_and_package(exported, ...) 内部就是把 ExportedProgram 喂给 Inductor 后端 + 生成 .so + 打包。整套链路与 JIT torch.compile 共享 90% 代码，差异在最后的 wrapper codegen（C++ vs Python）。

理解这条链路你就明白：v2.4+ 的 PyTorch 部署生态都建立在 Inductor 上。无论是训练时的 JIT 加速、还是部署时的 AOT 二进制，都是 Inductor 在底层支撑。Inductor 不只是一个编译器后端，是 PyTorch 通向”非 Python runtime 部署”的关键基础设施。

14.9.5.7 Inductor 在 LLM 推理上的关键优化路径

LLM 推理生态（vLLM / SGLang / TensorRT-LLM）大量用 torch.compile + Inductor。Inductor 给推理路径带来的核心收益：

• Decoder-only attention 的 fused kernel：QKV 投影 → split → SDPA → output projection 整段被 Inductor 识别 + fuse，相比 eager 加速 2-3x
• RMS Norm + Linear 的 fuse：Llama 类用 RMSNorm 替代 LayerNorm，Inductor 把 RMSNorm 与后续 Linear fuse 成单 kernel
• Rotary Embedding 与 attention 的协同：RoPE 计算与 attention QK score 在 Inductor 路径上能 fuse
• KV cache 写入与 attention 的 fuse：把 “compute K/V → 写到 cache → 用 cache 算 attention” 整段编进一个 kernel（vLLM 的 PagedAttention 用这条路）

这些优化都是”通用 fusion + pattern matcher 自动识别”的产物。不需要 LLM 框架开发者手写 SDPA / Norm / RoPE 的 fused kernel —— Inductor 自动生成。这是 v2.x 之后 LLM 推理引擎能快速演进的工程基础。

国内推理引擎（如 LMDeploy、TGI 中文化版）也大量基于这套路径。理解 Inductor 让你能在调优 LLM 推理时知道”哪些性能问题该让 Inductor 自动解决、哪些需要手写 CUDA”。

14.9.6 数学函数精度对齐

sin / cos / exp / log / sqrt / pow / erf 等数学函数有个常被忽略的工程问题：PyTorch 内置 ATen 实现与 Triton 库的数值精度可能不同。

举例：Triton 的 tl.exp 在某些 GPU 架构上用 __expf（fast intrinsic），精度比 expf 低一些。如果用户代码对精度敏感（如 numerical analysis 类训练），eager 与 compiled 的 loss 曲线可能分叉。

Inductor 的解法：

• 默认走 fast 数学函数（性能优先）
• 用户能通过 inductor.config.fast_math = False 强制走 IEEE-754 严格版
• 关键数学函数（如 exp 在 softmax 里）有专门的 codegen 路径与 ATen 对齐

这种”精度 vs 性能”取舍在 ML 训练里通常无影响（梯度有噪声，几个 ULP 差异不可见），但在 numerical 任务（如物理模拟、金融计算）必须知道这个开关存在。

14.9.6.5 几个用户视角的”为什么”

整章信息密集，最后用 Q&A 形式回答几个常被问的问题：

• 为什么 torch.compile 第一次特别慢？ —— 因为整条 Dynamo + AOTAutograd + Inductor + Triton 编译链路要跑一遍，几十个 Triton kernel 各编几秒。第二次起命中 FxGraphCache 直接 load
• 为什么我 fuse 不了？ —— 90% 是循环维度不对齐（一个 [N] 一个 [N, M]）或 reduction 隔断。开 TORCH_LOGS=fusion 看 scheduler 拒绝原因
• 为什么 Inductor 编出来的代码比我手写的 CUDA 快？ —— 不是单 kernel 比手写快，是 fusion 让多个手写 kernel 合并、整体内存读写减半
• 可以用 Inductor 编译我的 ResNet 部署吗？ —— 用 AOTI（§15.6.7）→ 输出 .pt2 → C++ 加载，整套不依赖 Python
• Inductor 与 TensorRT 比怎么样？ —— TensorRT 在静态推理场景峰值性能略高（专门优化 inference），Inductor 在训练 + 动态 shape + 易用性上完胜。生产推理服务取决于 latency / 部署灵活性的权衡

这些 Q&A 把整章串起来。每个问题的答案都对应前面某节的具体技术点，能反过来回查源码。

14.9.7 Scheduler 的 reorder passes

scheduler.py 还有一组 reorder pass（在 fusion 之后跑）：

• reorder_for_locality：把”读同一 buffer 的多个 op”挪到一起，让 cache 命中率高
• reorder_for_peak_memory：调整 op 顺序让 peak memory 最小
• reorder_communication_preserving_pin：分布式场景下让 collective 操作与 compute 充分 overlap

这些 reorder 在 fusion 之后做，因为 fusion 决定了哪些 op 已经合体（合体后内部顺序 fix），剩下的”独立 SchedulerNode 之间”还能调顺序。

reorder_for_peak_memory 与 §14.6.11 的 memory_planning 互补：memory_planning 是”已知 schedule 后做内存复用”，reorder_for_peak_memory 是”先调 schedule 让内存峰值更小”。两者一前一后让最终的内存使用接近最优。

14.9.8 Inductor 调试武器库

Inductor 编译失败 / 性能不达预期 / 数值不对时的诊断工具：

• TORCH_LOGS=inductor：开 Inductor 全套日志，看每段编译做了什么
• TORCH_LOGS=output_code：dump 生成的 Triton + wrapper Python 源码
• TORCH_LOGS=fusion：看 Scheduler 的 fusion 决策（哪些 op 合并、哪些拒绝、原因）
• TORCH_LOGS=schedule：看 reorder pass 后的最终调度顺序
• TORCH_LOGS=aot_graphs：看 AOTAutograd 输出给 Inductor 的 graph
• TORCH_COMPILE_DEBUG=1：把每次编译的中间产物（fx graph / lowered IR / Triton kernel）保存到 /tmp/torchinductor_debug/，方便事后分析
• torch._dynamo.config.verbose=True：Dynamo 的详细输出

这套日志组合让你能从”编译失败的 stack trace”一路追到”哪个算子的 lowering 拒绝了什么 fusion”。生产级 Inductor 调优必须熟悉这些工具。

特别有用的是 TORCH_COMPILE_DEBUG=1 产生的目录结构：

/tmp/torchinductor_debug/
  model_0/
    fx_graph_runnable.py     ← 可以独立运行的 fx graph repro
    fx_graph_readable.py     ← 人类友好的 fx graph
    fx_graph_transformed.py  ← 经过 fx_passes 后的 graph
    output_code.py           ← 最终 Triton kernel
    debug.log                ← 编译过程日志

把 fx_graph_runnable.py 单独跑能复现编译失败 —— 这是给 PyTorch 团队提交 Inductor bug PR 时的标准 repro 格式。

14.10 设计启示

Inductor 的几条核心思想：

第一：用更高级的 DSL 做 codegen 后端：选 Triton 而非 CUDA C++ 让 codegen 复杂度降一个数量级。这套思想可以借鉴到任何”自动生成 GPU 代码”的场景

第二：用 inner_fn 描述循环 body：把”逻辑循环”和”实际生成”解耦，让 fusion 决策能在纯函数式抽象上做，避免循环结构的干扰

第三：贪心 fusion + autotune 双保险：编译期用启发式决定融合范围（贪心可能不是全局最优），运行期 autotune 选 block size。两层一起把性能榨到接近最优

第四：ExternKernel 接现成库：不是所有算子都要自己生成。matmul / conv / attention 这些有几十年优化的算子用 cuBLAS / cuDNN / FlashAttention 调用即可，Inductor 只优化”算子之间的胶水”

下一章把 Dynamo / AOTAutograd / Inductor 串起来，加上 CUDA Graph，看 torch.compile 完整端到端流程。