Inductor

Active contributors: jansel, eellison, shunting314, desertfire, Chillee

Purpose

Inductor is the production compiler backend behind torch.compile. It takes an FX graph (typically produced by AOT Autograd) and lowers it to fast kernels: Triton on GPU and C++/OpenMP on CPU. It owns the IR, the scheduler, the codegen targets, and a battery of kernel templates for things like matmul, attention, and convolution.

The package is torch/_inductor/ and is one of the largest single packages in the repo.

Directory layout

Path	Contents
torch/_inductor/__init__.py	Top-level entry
torch/_inductor/compile_fx.py	The default BackendCompilerFunction that Dynamo calls
torch/_inductor/lowering.py	ATen op → Inductor IR lowering rules
torch/_inductor/decomposition.py	Inductor-specific decompositions
torch/_inductor/ir.py	The Inductor IR (Buffer, Loop, Pointwise, Reduction, ScatterMutation, …)
torch/_inductor/scheduler.py	Topological scheduling + fusion
torch/_inductor/codegen/	Code generators
torch/_inductor/codegen/triton.py	Triton GPU codegen
torch/_inductor/codegen/cpp.py	C++ CPU codegen
torch/_inductor/codegen/halide.py	Halide backend (experimental)
torch/_inductor/codegen/cuda/	CUDA C++ for templates (gemm/conv via CUTLASS)
torch/_inductor/codegen/rocm/	ROCm equivalents (Composable Kernel)
torch/_inductor/kernel/	Hand-written templates: mm, bmm, conv*, flex/, flash_attention.py
torch/_inductor/runtime/	Runtime support: triton heuristics, autotuning, hint extraction
torch/_inductor/fx_passes/	FX-level optimization passes (pre-grad and post-grad)
torch/_inductor/cudagraph_*.py	CUDA graph integration
torch/_inductor/freezing*.py	Constant folding / weight freezing
torch/_inductor/select_algorithm.py	Autotuning + algorithm selection
torch/csrc/inductor/	AOTInductor runtime + C-shim helpers

Key abstractions

Type	File	Purpose
GraphLowering	torch/_inductor/graph.py	Lowers an FX graph to Inductor IR
Buffer / ComputedBuffer	torch/_inductor/ir.py	A logical tensor in IR
Pointwise / Reduction	torch/_inductor/ir.py	Compute primitives
Scheduler	torch/_inductor/scheduler.py	Topo-orders nodes; performs fusion; emits kernels
TritonKernel	torch/_inductor/codegen/triton.py	A Triton kernel under construction
CppKernel	torch/_inductor/codegen/cpp.py	A C++ kernel under construction
select_algorithm.AlgorithmSelectorCache	torch/_inductor/select_algorithm.py	Caches autotune choices for matmul/conv templates

How it works

graph LR
    FX[FX graph from AOT autograd] -->|lowering.py| IR[Inductor IR]
    IR -->|decomposition.py| IR
    IR -->|fx_passes/| IR2[Optimized IR]
    IR2 -->|scheduler.py| Sch[Scheduler<br/>fuses pointwise + reductions]
    Sch -->|codegen/triton.py| Triton[Triton kernel src]
    Sch -->|codegen/cpp.py| Cpp[C++ kernel src]
    Triton -->|jit compile| Bin1[GPU binary]
    Cpp -->|cl/g++| Bin2[CPU binary]
    Bin1 & Bin2 --> Wrap[Compiled callable]

Lowering

compile_fx.compile_fx_inner is the entry point that AOT Autograd calls. It:

1. Runs decompositions (decomposition.py + the global torch/_decomp/) to break ATen ops into Inductor-friendly primitives.
2. Runs FX passes (fx_passes/) — these include reorder, fusion of small ops, group GEMM, normalization fusion, etc.
3. Calls GraphLowering.run to traverse the FX graph and call the lowering rule registered for each op (a function decorated with @register_lowering). The result is an Inductor IR DAG.

IR

Inductor IR is intentionally minimal. The main node kinds are:

• Pointwise — an elementwise expression over a numeric grid.
• Reduction — a reduction expression over a numeric grid.
• Scan / Sort / Gather — specialized loops.
• ExternKernel — a call into an external library (cuBLAS, cuDNN, custom kernels).
• TemplateBuffer — a hand-written kernel template (matmul, attention, convolution).

Each node knows how to compute its body lazily — bodies are sympy-flavored expressions over loop indices.

Scheduling and fusion

The Scheduler topologically orders the IR DAG into SchedulerNodes and fuses compatible neighbours:

• Pointwise neighbours that share an iteration space fuse into one kernel.
• Reductions can absorb a prologue of pointwise ops and an epilogue that runs over the reduced output.
• Memory-dependency analysis ensures correctness in the presence of mutations.

The scheduler is also where the loop-tiling heuristics for Triton kernels get decided.

Codegen

codegen/triton.py produces Python source for Triton kernels: @triton.jit decorated functions plus an outer Python wrapper that performs autotuning and kernel launch. Tile sizes (BLOCK_M, BLOCK_N, RBLOCK, etc.) are chosen by configs in runtime/triton_heuristics.py.

codegen/cpp.py produces C++ source: vectorized loops that go through aten/src/ATen/cpu/vec/ SIMD primitives. Output is compiled with g++/clang through a small torch/_inductor/cpp_builder.py that handles caching.

Templates and autotuning

For matmul, batched matmul, attention, and convolution Inductor uses templates (kernel/mm.py, kernel/flash_attention.py, kernel/flex/) that pick between several algorithm choices (different tile sizes, ATen bmm, cublas, cuDNN, CUTLASS). select_algorithm.py runs a quick benchmark of each candidate on the actual input shapes and caches the winner.

CUDA Graphs

cudagraph_trees.py and cudagraph_utils.py integrate CUDA Graphs into the runtime so that consecutive compiled regions can be replayed without per-launch CPU overhead. Trees of CUDA graphs share their input/output memory pools.

AOTInductor

For deployment scenarios (no Python at runtime), torch._inductor ships an "ahead-of-time" mode that compiles the graph into a .so plus a small C runtime in torch/csrc/inductor/aoti_runtime/. Used by ExecuTorch, libtorch deployment, and Meta's internal inference stack.

Integration points

• AOT Autograd is the upstream producer of FX graphs. See AOT Autograd.
• torch.compile is the user-facing entry point that pulls everything together. See Features / torch.compile.
• Triton is a hard runtime dep for GPU codegen; the pinned commit lives in .ci/docker/ci_commit_pins/triton.txt.
• AOTInductor is the C runtime side at torch/csrc/inductor/.

Entry points for modification

• New ATen op support → add a @register_lowering rule in lowering.py and (optionally) a decomposition.
• New codegen target → add a directory under codegen/ and a Backend registration; see codegen/halide.py as the smallest example.
• New fusion pattern → an fx_passes/ pass.
• New kernel template → add a file under kernel/ and register algorithm choices in select_algorithm.py.
• For debugging: TORCH_COMPILE_DEBUG=1 dumps every IR / scheduler / codegen artefact to torch_compile_debug/.

Key source files

File	Purpose
torch/_inductor/compile_fx.py	Entry point, top-level orchestration
torch/_inductor/graph.py	GraphLowering
torch/_inductor/lowering.py	ATen op → IR rules
torch/_inductor/ir.py	IR node definitions
torch/_inductor/scheduler.py	Scheduling + fusion
torch/_inductor/codegen/triton.py	Triton codegen
torch/_inductor/codegen/cpp.py	C++ codegen
torch/_inductor/select_algorithm.py	Autotuning
torch/_inductor/kernel/mm.py	Matmul template
torch/_inductor/kernel/flex/	FlexAttention
torch/_inductor/cudagraph_trees.py	CUDA Graph integration
torch/_inductor/runtime/triton_heuristics.py	Tile-size heuristics
torch/csrc/inductor/aoti_runtime/	AOTInductor C runtime