Inside PyTorch(2) — How PyTorch Compiles for GPUs
🔍 Does TorchInductor Build CUDA Graphs from Triton Kernels?
Broadly, yes. TorchInductor optimizes the PyTorch computation graph and commonly generates GPU kernels through Triton. CUDA Graphs can then capture those kernels together with other CUDA operations and replay the prepared workload with lower CPU launch overhead.
① GPU Kernels and CPU Launches
A GPU kernel is a function executed in parallel by many GPU threads. For example, the following PyTorch code contains three tensor operations.
1
2
3
x = a + b
y = torch.relu(x)
z = y * c
Without compiler fusion, these operations may result in multiple GPU kernel launches.
1
2
3
4
5
add kernel
↓
relu kernel
↓
multiply kernel
The CPU does not perform these calculations itself. Instead, host code uses the CUDA Runtime and Driver to submit work to a CUDA stream.
1
2
3
4
5
6
7
CPU
↓
CUDA Runtime / Driver
↓
CUDA Stream
↓
GPU Kernel
A kernel launch is normally asynchronous. The CPU submits the work, and the GPU executes it when its dependencies are satisfied. This means GPU execution has two different costs.
- Kernel execution time: the time spent performing the calculation on the GPU.
- Kernel launch overhead: the CPU-side cost of preparing and submitting the kernel.
When many short-lived kernels are launched, CPU launch overhead can become a bottleneck even if each kernel performs only a small amount of work.
② Why Did Triton and TorchInductor Appear?
Writing efficient GPU kernels directly in CUDA C++ requires detailed control over parallelism, memory access, synchronization, and hardware-specific behavior.
Triton provides a different programming model. It is a Python-based language and compiler for writing parallel GPU programs. Developers describe work at a higher level than CUDA threads while still controlling important details such as memory loads, reductions, tiling, and program instances. A simplified Triton kernel looks like this.
1
2
3
4
5
6
7
8
9
10
11
12
13
14
@triton.jit
def add_kernel(
x_ptr,
y_ptr,
output_ptr,
n_elements: tl.constexpr,
):
offsets = tl.program_id(0) * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
mask = offsets < n_elements
x = tl.load(x_ptr + offsets, mask=mask)
y = tl.load(y_ptr + offsets, mask=mask)
tl.store(output_ptr + offsets, x + y, mask=mask)
Triton makes individual GPU kernels easier to express, but it does not decide how an entire PyTorch model should be optimized.
That is the role of TorchInductor. PyTorch Eager Mode dispatches each tensor operation for execution as soon as it is encountered, without first optimizing the surrounding computation graph.
1
2
3
4
5
6
7
Python Program
↓
PyTorch Operation
↓
Execute Immediately
↓
ATen / CUDA Kernel
Because each operation is handled separately, Eager Mode has limited visibility into the surrounding computation. torch.compile() changes this execution path.
1
2
3
4
5
6
7
8
9
PyTorch Program
↓
TorchDynamo
↓
AOTAutograd
↓
TorchInductor
↓
Optimized CPU or GPU Code (Triton Kernel or ATen / cuBLAS / cuDNN)
TorchDynamo captures compilable regions of the Python program. AOTAutograd prepares forward and backward graphs. TorchInductor then analyzes those graphs and generates optimized code for the target device.
③ TorchInductor vs Triton
TorchInductor and Triton are both compilers, but they operate at different levels.
| TorchInductor | Triton |
|---|---|
| Receives PyTorch computation graphs | Receives Triton kernel programs |
| Examines multiple tensor operations | Implements an individual GPU kernel |
| Decides fusion and scheduling | Compiles the kernel for GPU execution |
| Plans intermediate buffers and memory reuse | Expresses loads, stores, reductions, and parallel work |
| May generate Triton code or use external libraries | Does not optimize the entire PyTorch model |
Consider the following expression.
1
output = torch.relu((a + b) * c)
TorchInductor may determine that the addition, multiplication, and activation can be fused.
1
2
3
4
5
FX Graph
│
├── add
├── multiply
└── relu
Instead of launching three independent kernels, TorchInductor can generate one fused Triton kernel.
1
triton_fused_add_mul_relu
The resulting execution changes from this:
1
2
3
launch add
launch multiply
launch relu
to this:
1
launch fused_add_mul_relu
Fusion can reduce:
- the number of kernel launches,
- intermediate tensor allocations,
- repeated global-memory reads and writes.
However, TorchInductor does not generate every GPU operation as a Triton kernel.
Operations such as matrix multiplication or convolution may use optimized implementations from ATen, cuBLAS, cuDNN, or compiler-generated templates. A compiled region can therefore contain a mixture of Triton kernels and external library calls.
The relationship is better represented as follows.
1
2
3
4
5
6
7
8
9
10
11
12
13
PyTorch FX Graph
↓
TorchInductor
│
├── Fusion
├── Scheduling
├── Memory Planning
└── Code Generation
│
┌───────┴────────┐
↓ ↓
Triton Kernels External Libraries
ATen / cuBLAS / cuDNN
TorchInductor determines what code should be generated or called.Triton helps implement and compile GPU kernels selected by that process.
④ Why Are CUDA Graphs Still Needed?
TorchInductor and Triton reduce the amount of GPU work by applying optimizations such as kernel fusion. They do not necessarily eliminate CPU launch overhead. A compiled model may still execute several GPU operations.
1
2
3
4
5
6
7
Triton fused kernel
↓
cuBLAS matrix multiplication
↓
Triton reduction kernel
↓
ATen CUDA kernel
Without CUDA Graphs, the CPU submits these operations separately during every iteration.
1
2
3
4
5
6
7
8
9
10
11
12
13
14
Iteration 1:
launch A
launch B
launch C
Iteration 2:
launch A
launch B
launch C
Iteration 3:
launch A
launch B
launch C
Even when the same sequence repeats, the CPU and CUDA Driver must perform launch-related work again. CUDA Graphs introduce a different work-submission model. A graph describes CUDA operations and their dependencies separately from execution. After the graph has been prepared, it can be launched repeatedly. The lifecycle can be simplified into three steps.
Definition or Capture
CUDA records the operations and their dependencies.
1
2
3
4
5
6
kernel A
↓
kernel B
├──────────┐
↓ ↓
kernel C memcpy
Instantiation
The graph definition is converted into an executable graph.
1
2
3
CUDA Graph
↓
Executable CUDA Graph
Replay
The executable graph is launched repeatedly.
1
2
3
Iteration 1: graph launch
Iteration 2: graph launch
Iteration 3: graph launch
The main improvement is on the CPU side.
Without CUDA Graphs
1
2
3
4
5
6
Every iteration:
CPU:
- prepare & launch kernel A
- prepare & launch kernel B
- prepare & launch kernel C
With CUDA Graphs
1
2
3
4
5
6
First execution:
- capture CUDA workload
- instantiate executable graph
Every iteration:
- launch executable graph
CUDA Graphs do not necessarily make the code inside an individual Triton kernel faster. They reduce the cost of repeatedly submitting the overall GPU workload.
⑤ What Can a CUDA Graph Contain?
A CUDA Graph is not specific to Triton. It can represent multiple types of CUDA operations, including:
- Triton-generated kernels,
- ATen CUDA kernels,
- cuBLAS or cuDNN operations,
- memory copies,
- memory-set operations,
- CUDA events,
- child graphs.
Therefore, the following description is incomplete.
Triton kernels are converted into a CUDA Graph.
A more accurate description is:
CUDA Graphs capture a GPU workload that may contain Triton kernels together with other CUDA operations.
The graph represents both the operations and their dependencies. It is not merely a text list of kernel names.
⑥ How TorchInductor, Triton, and CUDA Graphs Work Together
The complete execution path can be summarized as follows.
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
PyTorch Program
│
▼
TorchDynamo
Captures a computation graph
│
▼
AOTAutograd
Prepares forward and backward graphs
│
▼
TorchInductor
Fuses and schedules operations
│
├──────────────────┐
▼ ▼
Triton Kernels Library Operations
ATen / cuBLAS / cuDNN
│ │
└────────┬─────────┘
▼
CUDA Operations
│
▼
CUDA Graph Capture
│
▼
Executable CUDA Graph
│
▼
Graph Replay
Each component optimizes a different layer.
| Component | Optimization Layer |
|---|---|
| TorchInductor | Computation graph and generated program |
| Triton | Individual GPU kernel implementation |
| CUDA Graph | CPU-to-GPU work submission |
TorchInductor decides how the computation should be organized. Triton generates GPU kernels for compatible regions. CUDA Graphs reduce the overhead of repeatedly launching the resulting GPU workload.
⑦ What Changed After CUDA Graphs?
Suppose one model iteration contains four short GPU operations.
Conventional execution
1
2
3
4
CPU → launch kernel A
CPU → launch kernel B
CPU → launch kernel C
CPU → launch kernel D
The launch path is repeated for every iteration.
CUDA Graph execution
1
CPU → launch executable graph
The GPU still performs kernels A, B, C, and D. The difference is how the work reaches the GPU.
| Before CUDA Graphs | With CUDA Graphs |
|---|---|
| Operations are submitted individually | A prepared workload is submitted as a graph |
| CPU launch work repeats for every operation | Much of the preparation is reused |
| Flexible for changing execution paths | Best suited to stable, repeated workloads |
| Launch overhead grows with many short operations | CPU overhead can be reduced significantly |
| Ordinary CUDA memory allocation behavior | Graph execution may require persistent workspace memory |
CUDA Graphs are most effective for workloads that repeatedly execute the same sequence of GPU operations with stable tensor shapes and execution paths, especially when many short-lived kernels make CPU launch overhead a bottleneck. They are less suitable for workloads with frequently changing tensor shapes or control flow, unsupported operations, or memory layouts that prevent stable address reuse.
PyTorch exposes this optimization through torch.compile() modes and options.
1
2
3
4
compiled_model = torch.compile(
model,
mode="reduce-overhead",
)
The reduce-overhead mode uses CUDA Graphs when the workload is compatible. CUDA Graphs can also be requested through a TorchInductor option.
1
2
3
4
compiled_model = torch.compile(
model,
options={"triton.cudagraphs": True},
)
Enabling the option does not guarantee that every compiled region will be captured. PyTorch may skip CUDA Graphs when the workload violates capture requirements.
Key Takeaway
TorchInductor, Triton, and CUDA Graphs do not perform the same optimization.
| TorchInductor | Triton | CUDA Graph |
|---|---|---|
| Optimizes the PyTorch computation graph | Generates and compiles GPU kernels | Captures and replays CUDA workloads |
| Decides fusion and scheduling | Implements kernel-level parallel computation | Reduces repeated CPU launch overhead |
| Operates above individual kernels | Operates at the kernel level | Operates at the execution-submission level |
CUDA Graphs do not replace Triton. Triton does not replace TorchInductor. They optimize different layers of the same execution path.
TL;DR
Q. What does TorchInductor use to generate GPU kernels?
TorchInductor commonly uses Triton as a key GPU code-generation component. It can also call ATen kernels, compiler templates, and optimized libraries such as cuBLAS or cuDNN.
Q. Are TorchInductor and Triton the same compiler?
No. TorchInductor optimizes a PyTorch computation graph, while Triton implements and compiles individual GPU kernels.
Q. What does Triton optimize?
Triton expresses kernel-level parallel computation, including memory access, tiling, reductions, and GPU program instances.
Q. Does generating Triton kernels remove CPU overhead?
Not completely. Kernel fusion can reduce the number of launches, but the CPU must still submit the remaining GPU operations.
Q. What problem do CUDA Graphs solve?
CUDA Graphs reduce CPU launch overhead by defining or capturing a repeated CUDA workload once and launching the prepared executable graph multiple times.
Q. Do CUDA Graphs make individual Triton kernels faster?
Not necessarily. The kernel normally performs the same GPU calculation. CUDA Graphs optimize how the overall workload is submitted.
Q. Do CUDA Graphs contain only Triton kernels?
No. A graph may contain Triton kernels, ATen kernels, library operations, memory copies, events, and other CUDA operations.
Q. Does inductor_cudagraphs mean that CUDA Graphs contain kernels generated by Triton?
Broadly, yes. A more complete answer is: TorchInductor generates or selects GPU operations, including Triton kernels, and PyTorch can capture those operations into a CUDA Graph for repeated execution. Triton does not create the CUDA Graph itself, and the graph is not limited to Triton kernels.