필사 모드: GPU Compiler and Framework Landscape — One Problem, Turning a Graph into a Kernel, a Different Answer at Every Layer

• Introduction — Nobody Wrote the Kernel the Profiler Pointed To
• The One Shared Problem — Turning a Graph into a Kernel
• The Bottom Layer — NVCC, PTX, SASS, and LLVM's GPU Backends
• MLIR and What's Built on It — A Framework for Building IRs, and Triton
  • Built on Top of MLIR — The Path Triton Takes to Build a Kernel
• What torch.compile and Inductor Actually Do
• XLA, IREE, and TVM — Taking the Whole Graph
• Which Layer to Know, and When
• Closing — Layers Split Along Authority, Not Performance
• References

Introduction — Nobody Wrote the Kernel the Profiler Pointed To

Profile a training step with Nsight Compute and you get the name of the kernel eating the most time. The trouble starts when that name looks like triton_poi_fused_add_mul_native_layer_norm_7. Searching the repository turns up nothing, because nobody wrote it. A machine did.

Figuring out what actually produced it means answering questions like these. At which stage was the kernel generated? Did PyTorch produce it, did Triton, or was it the LLVM layer underneath? If you want to rename it — or really, if you want to change the code — which layer do you actually touch?

The question is hard not because there are so many layers, but because each layer has a different name, belongs to a different organization, and keeps its own documentation, so they all look unrelated to each other. NVCC, LLVM, MLIR, Triton, Inductor, XLA, IREE, and TVM look like they belong to entirely different worlds.

In reality, they all solve the same problem. This post states that problem first, then lays out which slice of it each layer has claimed. If the earlier posts in this series covered writing and modifying kernels by hand, three layers of kernel authoring, tuning an inference engine, and how the two vendor stacks differ, this post is the coordinate system they all sit on.

Versions and facts were checked directly against each project's official documentation and release notes on August 2, 2026. This landscape can shift within 6 months, so the text explicitly distinguishes what was verified from what wasn't.

The One Shared Problem — Turning a Graph into a Kernel

Whatever they're called, every tool in this post shares the same input and output.

The input is a computation graph — nodes like matrix multiplication, addition, activation functions, and normalization, connected by tensors. The output is machine code that actually runs on the GPU. And what has to happen in between is strikingly similar across all of them.

• A fusion decision. Which nodes get merged into a single kernel. Merging means intermediate tensors don't have to be written to global memory, which saves memory bandwidth. The XLA documentation calls fusion "the single most important optimization" in XLA.
• A layout decision. What order tensors get laid out in memory. The same operation can end up coalesced or not, depending purely on layout.
• A tiling decision. What size chunks a large matrix gets cut into and handed out to each thread block. Chunks too large and you run out of shared memory and registers; too small and reuse drops.
• A choice between calling a library and generating code. Is this operation faster as a call into cuBLAS, or faster as a hand-generated kernel?
• Lowering. The stage that turns all those decisions into actual instructions.

Multiple layers exist because these decisions differ in nature. Fusion can only be judged by looking at the whole graph; register allocation can only be judged at the instruction level. It's hard for a single abstraction to do both well. So the upper layers look at the graph, the lower layers look at instructions, and middle layers sit in between.

Summarized, it looks like this.

Keep this table in mind while reading on, and every tool below will slot into a clear position.

The Bottom Layer — NVCC, PTX, SASS, and LLVM's GPU Backends

Let's start from the very bottom. This is the one place whose structure hasn't fundamentally changed in 20 years.

Feed a .cu file into nvcc and it splits into host code and device code. Device code comes out in two forms: PTX and CUBIN. The CUDA 13.3 compiler documentation describes the procedure like this: device functions are compiled into PTX or CUBIN, those get packed into a fatbinary container, and the original source is transformed into standard C++ carrying that fatbinary, which is then handed to the host compiler. At run time, "the CUDA runtime inspects the embedded fatbinary" to pick the image that matches the current GPU.

The concept to understand here is the distinction between the virtual architecture and the real architecture.

# compute_90 = virtual architecture, the instruction-set generation PTX is generated for
# sm_90      = real architecture, the binary for the actual chip
nvcc -gencode arch=compute_90,code=sm_90 \
     -gencode arch=compute_90,code=compute_90 \
     kernel.cu -o kernel

# The first line puts SASS for Hopper into the fatbinary; the second line puts PTX itself in.
# The second line is what lets the driver JIT-revive it on future GPU generations.

As the documentation puts it, when PTX is bundled in, "the CUDA runtime dynamically compiles the embedded PTX code" whenever there's no binary load image matching the current GPU. That's the entirety of forward compatibility. And that's why PTX doesn't go away. PTX looks like human-readable assembly, but it's actually a distribution format and a stable contract. SASS's encoding changes with every generation and is undocumented, while PTX is documented and preserves backward compatibility.

There's a point in practice where this distinction actually bites. According to the PyTorch 2.13 release notes, CUDA 13 builds no longer bundle ptxas into the binary. That means the tool that turns PTX into SASS isn't in the wheel, and environments that depend on the PTX JIT path, or custom build scripts that call ptxas directly, can suddenly break. Not needing to know about a layer is not the same as not needing that layer to exist.

ptxas handles the final leg from PTX to SASS. Register allocation and instruction scheduling happen at this stage, and the register count that governs occupancy is decided here too. When performance shifts after nothing but a CUDA toolkit minor-version bump — kernel source untouched — it's almost always because ptxas scheduling changed underneath you. You can check this with the tools below.

# See the SASS that actually got generated
cuobjdump -sass ./kernel | head -40

# Register and shared memory usage (inputs to the occupancy calculation)
nvcc -Xptxas -v -arch=sm_90 -c kernel.cu
# ptxas info : Used 64 registers, 8192 bytes smem, 384 bytes cmem[0]

# Extract just the PTX
nvcc -ptx -arch=compute_90 kernel.cu -o kernel.ptx

But CUDA C++ isn't the only thing that produces GPU machine code. LLVM has a backend called NVPTX that lowers LLVM IR into PTX. Triton, XLA, and nearly every other code generator eventually passes through this same backend. In other words, PTX is simultaneously a CUDA artifact and an LLVM artifact.

AMD's side has the same structure under a different name. LLVM's AMDGPU backend lowers LLVM IR into AMDGCN assembly, which becomes the code object. AMD, however, doesn't use upstream LLVM as-is — it maintains a fork. The ROCm documentation states plainly that its compiler is "a fork of llvm/llvm-project," names amdclang++ as the default compiler, and describes hipcc as a driver that "invokes clang or nvcc, passing the appropriate include and library options." At the time this was checked, the documentation listed the llvm-project version as 22.0.0.

Maintaining a fork isn't a bad sign in itself — new chip support and features are needed faster than the upstream release cycle can deliver them, and a good portion of AMD's changes do land upstream eventually. But it matters in practice. Third-party compiler projects targeting AMD GPUs have to reconcile their own LLVM version against ROCm's LLVM version, and that's one strand of the porting friction covered in the post comparing the two vendor stacks.

MLIR and What's Built on It — A Framework for Building IRs, and Triton

From here, the landscape changes.

LLVM IR was designed for CPUs — scalars and vectors, basic blocks, SSA. There's no concept of a tensor, no concept of nested loops. Express a matrix multiply in LLVM IR and it's already low-level code with loops unrolled; none of the information needed to decide something like "let's tile this loop" survives.

So every ML compiler in the late 2010s ended up repeating the same work: define a new tensor-level IR, write a new pass infrastructure from scratch, build a new parser, verifier, and printer, and finally write new code to lower it down to LLVM IR. XLA did it, TVM did it, and every hardware vendor did it on their own.

MLIR's proposal is to eliminate that repetition. Instead of defining one more IR, it provides infrastructure for defining IRs. A bundle of operations and types is called a dialect, and each project defines its own dialect while sharing the pass management, verification, serialization, and lowering framework.

The key point is that dialects coexist. A single module can mix operations at different levels of abstraction, and lowering becomes the process of gradually resolving that mixture.

// Two dialects at different levels of abstraction coexist in the same function.
// linalg preserves the meaning "this is a matrix multiply,"
// while scf/arith have already been lowered down to loops and scalar operations.
func.func @mm(%A: tensor<128x256xf32>, %B: tensor<256x64xf32>)
    -> tensor<128x64xf32> {
  %init = tensor.empty() : tensor<128x64xf32>
  %C = linalg.matmul
      ins(%A, %B : tensor<128x256xf32>, tensor<256x64xf32>)
      outs(%init : tensor<128x64xf32>) -> tensor<128x64xf32>
  return %C : tensor<128x64xf32>
}

A single line of linalg.matmul carries the information "this is a matrix multiply" intact. A tiling pass can look at that information and cut loops accordingly. Lower the same thing down to LLVM IR first, and that information is gone — reconstructing loops afterward becomes far harder. The later you discard information, the better decisions you can make — that's MLIR's reason for existing, compressed into one sentence.

MLIR doesn't compile anything on its own. What gets built on top of it is everything. And a lot has been built on top. Triton's IR, IREE's IR, parts of XLA, and various hardware vendors' in-house compilers all sit on MLIR. The fact that projects with completely different names share the same foundation matters for understanding this landscape.

Built on Top of MLIR — The Path Triton Takes to Build a Kernel

Triton is both a kernel-authoring language and an MLIR-based compiler. The post comparing three layers of kernel authoring already covered Triton as a language, so here we look at its position as a compiler.

The pipeline is clearly defined. The Triton plugin extension documentation the PyTorch team published in July 2026 lays out the stages like this.

Python kernel function
      ↓  (AST traversal)
TTIR    — Triton IR. Tile-level operations. Hardware-agnostic.
      ↓  (layout decisions, thread mapping)
TTGIR   — TritonGPU IR. Warp placement and shared memory are determined.
      ↓
LLVM IR
      ↓
PTX (NVIDIA)  or  AMDGCN (AMD)

These four stages compress the "layer" structure laid out earlier in this post into one picture. TTIR is the kernel layer, TTGIR is the boundary between the kernel and instruction layers, and everything below LLVM IR is the instruction layer. Users write code at roughly the TTIR level and leave the rest to the compiler. That's the deal Triton offers.

According to the repository README, the supported backends are NVIDIA (Compute Capability 8.0 and above) and AMD (ROCm 6.2 and above), with CPU marked as "in development." The CPU backend has several efforts underway, but as of when this was checked it remains experimental — not something to rely on in practice yet.

The change worth watching at this layer in 2026 is the plugin extension system. It shipped in pytorch-triton 3.7, which PyTorch distributes, and PyTorch 2.13 pins Triton at 3.7.1. Previously, adding a custom pass or dialect meant forking Triton. A fork quickly falls behind upstream, merge conflicts pile up, and you end up stuck on a stale release that misses new hardware support. The plugin system lets you load a shared library at runtime to insert a pass at an arbitrary point in the pipeline, disable a specific pass, or replace an entire stage.

# Plugins are loaded via an environment variable. Import order matters.
import os, sysconfig

dist_packages = sysconfig.get_paths()["purelib"]
os.environ["TRITON_PLUGIN_PATHS"] = os.path.join(
    dist_packages, "utlx_plugin", "libutlx.so"
)

import triton
import triton.language as tl
import utlx_plugin as tlx      # TLX extension ops come in here

The first consumer is Meta's TLX (Triton Language Extensions). Operations like tlx.local_alloc, tlx.async_load, and tlx.async_dot let you handle shared memory buffers and asynchronous pipelining explicitly, inside the kernel. In other words, it's a move to hand some of the decisions Triton used to make automatically back to the person writing the kernel. The published numbers show it ahead of cuBLAS by 2 to 4 percent on H100 persistent GEMM, and ahead of rocBLAS by 12 to 15 percent on MI350 pipelined GEMM. The point is that a kernel faster than the vendor library is coming out of a compiler language.

There's a pattern visible here. Abstraction layers don't only move upward. Where Triton moved up by hiding threads, TLX moves back down by re-exposing asynchronous instructions. The height of an abstraction is determined not by performance, but by which decisions a human gets to hold onto.

What torch.compile and Inductor Actually Do

This is the layer PyTorch users run into most often and understand least. Behind a single line of torch.compile(model), three stages run in sequence.

TorchDynamo intercepts Python bytecode and pulls out an FX graph. Pure tensor operations go into the graph; things like Python list manipulation or a print call get pushed out of it. That point is a graph break. When the graph gets cut into multiple pieces, you pay a cost dropping back into Python and out again around each piece, and small pieces also lose fusion opportunities. Most cases where torch.compile doesn't speed things up as much as expected aren't because the kernels are bad — it's because the graph got cut up.

AOTAutograd splits that graph into two ATen-level graphs, forward and backward, and decides what gets saved versus what gets recomputed.

TorchInductor is what actually emits code. For a GPU target it generates Triton kernels; for a CPU target it generates C++ and OpenMP. That's where names like the triton_poi_fused_... we saw earlier come from. The name itself carries information: poi stands for pointwise, and whatever follows fused_ is the list of operations that got merged together.

The generated code can be inspected directly. That's the fastest way to understand this layer.

import os
os.environ["TORCH_COMPILE_DEBUG"] = "1"        # Leaves the generated code on disk
os.environ["TORCH_LOGS"] = "output_code,graph_breaks"

import torch

def block(x, w, b):
    return torch.nn.functional.gelu(x @ w + b)

x = torch.randn(1024, 512, device="cuda", dtype=torch.bfloat16)
w = torch.randn(512, 512, device="cuda", dtype=torch.bfloat16)
b = torch.randn(512, device="cuda", dtype=torch.bfloat16)

compiled = torch.compile(block)
out = compiled(x, w, b)

# A way to check graph breaks on their own
explanation = torch._dynamo.explain(block)(x, w, b)
print("Graph count:", explanation.graph_count)
print("Break count:", explanation.graph_break_count)

Turn on TORCH_LOGS="output_code" and the Triton source Inductor generated prints to standard output. You'll typically see the matrix multiply left as a cuBLAS call, with GELU and the bias addition merged into a single pointwise kernel. That kernel is exactly the identity of the nameless kernel the profiler pointed to earlier.

One important change happened at this layer in 2026. Starting with PyTorch 2.13, Inductor gained an additional path besides Triton: the CuTeDSL path. It's a DSL built on top of NVIDIA CuTe, and according to the release notes it's a second high-performance path targeting transformer GEMM and RMSNorm, aimed at "generating better-quality matrix-multiply code even without Triton." The same release also moved kernel compilation from a thread pool to a subprocess pool, because Python's GIL had been a bottleneck on compile time.

The same release introduced torch.compiler.set_default_backend. It's an API that sets the process-wide default backend the same way torch.set_default_dtype does, so you no longer need to attach backend= to every torch.compile call. It's a meaningful change for hardware vendors building out-of-tree backends.

To sum up, torch.compile is a graph compiler that delegates kernel generation to the layer below. That layer used to be Triton alone; now it's two.

XLA, IREE, and TVM — Taking the Whole Graph

If Inductor is a graph compiler attached to PyTorch, there's a separate lineage that aims to take in a graph independent of any one framework.

XLA is the oldest and most widely used. It's JAX's only execution path, TPU's only execution path, and PyTorch connects to it too, via PyTorch/XLA. Its input is HLO, and the contract between framework and compiler is handled by StableHLO. StableHLO is an operation set that promises 5 years of backward compatibility and 2 years of forward compatibility, which is exactly what makes it possible to open a graph saved years ago in a compiler built years later.

XLA:GPU's code generation approach lines up exactly with this post's subject. In the official documentation's own words, XLA:GPU uses "a combination of native (PTX-via-LLVM) emitters and a TritonIR emitter." This splits into three branches.

1. Common operations just call cuBLAS, cuDNN, or NCCL directly.
2. Operations with a recognizable pattern, like reductions or transposes, get LLVM IR generated directly and lowered to PTX.
3. Advanced fusions involving matrix multiplies or softmax convert the HLO fusion into TritonIR, pick tile parameters, then call into Triton to get PTX back.

The third is the important one. XLA and Triton aren't competitors — they're in a calling relationship. XLA makes the graph-layer decisions and hands the kernel-layer decisions off to Triton. It's essentially the layer split from the earlier table, implemented directly. Finally, the XLA runtime moves the sequence of kernel calls and library calls into its own MLIR dialect called RuntimeIR, and extracts a CUDA graph from there.

IREE solves the same problem across a much wider range of targets. In the README's words, it's an MLIR-based compiler and runtime that "lowers ML models to a unified IR that scales up to meet the needs of data centers and scales down to meet the constraints of mobile and edge deployments." The repository's topic tags list JAX, PyTorch, ONNX, and TensorFlow as frontends, and CUDA, ROCm, Vulkan, and SPIR-V as targets. It joined the LF AI & Data Foundation as a sandbox-stage project in May 2024.

IREE's distinguishing trait is that it designs the compiler and runtime together. Its output is a self-contained executable module rather than an object that needs a Python process, so the same pipeline reaches all the way to embedded targets with no Python at all. Running into IREE in a data center is still rare; as of when this was checked, the largest adoption case was AMD pushing it as part of its own stack.

TVM is the project whose position in this landscape has shifted the most. Around 2018 it was the only serious open-source deep learning compiler, and it drew attention with the claim that autotuning could beat vendor libraries. Today it's a cross-layer design using Relax as its graph-level representation and TensorIR as its tensor-level representation, and its documented design goal is "making most transformations customizable from Python, to make ML compilers accessible."

To put it plainly, you almost never run into TVM in frontier LLM training and serving anymore. Inductor and XLA took that spot. Where TVM is still alive is elsewhere: hardware with no vendor library, edge targets that can't run a Python runtime, and deployment paths like MLC LLM that run LLMs in the browser or on mobile. The original claim that autotuning beats a tuned library is hard to sustain where cuBLAS and CUTLASS exist, and still holds up where they don't.

Compared in one table, it looks like this.

Which Layer to Know, and When

Having many layers doesn't mean you need to know all of them. You only need to drop down when it's actually necessary. What signals that "necessary" moment is the most practical part of this post.

In my experience, 90 percent of practitioners never get past the top two rows. Eliminating a single graph break usually pays off more than hand-optimizing a kernel. The lower you go, the bigger the potential payoff, but the time it costs climbs far more steeply.

If you want to look at the Triton IR directly, here's how to dump it.

# Leaves each stage's IR behind as a file
export TRITON_KERNEL_DUMP=1
export TRITON_DUMP_DIR=/tmp/triton_ir
python train.py

ls /tmp/triton_ir/*/
# *.ttir     TTIR   — tile-level, hardware-agnostic
# *.ttgir    TTGIR  — layout and warp placement decided
# *.llir     LLVM IR
# *.ptx      PTX
# *.cubin    final binary containing SASS

Opening the same kernel's five files side by side gets you to understanding faster than reading this entire post did. Working from top to bottom, you can watch exactly what gets decided and what disappears at each step.

Finally, here's what this research could not verify.

• Triton's stable release version and date. The repository README doesn't state a version, and the only figure confirmed was the 3.7.1 that PyTorch pins. Triton's own release numbering and the pytorch-triton numbering PyTorch distributes may not match.
• IREE's latest release tag and date. It wasn't shown in the README and wasn't checked separately.
• Apache TVM's recent release timing and activity level. The design direction was confirmed from project documentation, but the release history itself wasn't looked up directly. The description above is a judgment about adoption, not a claim backed by statistics.
• Helion. It was confirmed that it's a kernel DSL PyTorch has been pushing in 2026 and that its TPU backend compiles to Pallas, but its relationship to Triton on the NVIDIA and AMD paths could not be confirmed from the documentation available.

Closing — Layers Split Along Authority, Not Performance

Folded down to one sentence, this post's map reads like this: layers don't split along speed, they split along who gets to hold which decision.

The graph layer takes fusion and layout. In exchange, it gives up the details of any one kernel. The kernel layer takes tiling and pipelining. In exchange, it can't see the whole graph. The instruction layer takes registers and scheduling. In exchange, it doesn't know what's actually being computed. That's why no layer can replace another, and it's the same structure CPU compilers split into 30 years ago.

So there's a fixed question to ask whenever a new project name shows up. Not what's new about it, but which decision it took, and from whom. Triton took thread mapping away from humans. TLX took asynchronous instructions back from the compiler. Inductor took fusion away from the user. MLIR took IR infrastructure away from individual projects.

That one question is usually enough to place a new framework somewhere on this map. And that's knowledge that lasts longer than chasing every name that changes every 6 months.

References

• NVIDIA CUDA Compiler Driver NVCC 13.3 — compilation trajectory and virtual vs. real architectures
• ROCm LLVM Project documentation — amdclang, hipcc, comgr
• Triton repository README — supported backends and the MLIR-based rewrite
• Triton Plugin Extensions: TLX and Custom Compiler Passes — PyTorch blog, 2026-07-15
• PyTorch 2.13 release notes — CuTeDSL backend, Triton 3.7.1 pin, CUDA 13 changes
• XLA:GPU Architecture Overview — native emitters and the TritonIR emitter
• StableHLO — the compatibility layer between frameworks and compilers
• IREE repository README — an MLIR-based compiler and runtime
• Apache TVM — the cross-layer design of Relax and TensorIR
• MLIR project — dialects and progressive lowering