Exploring how MLIR works: the compiler rewiring the AI stack
From tensor graphs to machine code, MLIR is quietly becoming the abstraction layer connecting modern AI frameworks to increasingly specialized hardware.
Qualcomm just agreed to pay roughly $3.9 billion for a 150 person compiler company. Tesla says it rebuilt the FSD compiler and runtime on the same infrastructure. NVIDIA’s newest Python kernel DSL compiles through it. That infrastructure is MLIR.
This is the full tour: the object model, the dialect system, one MLP block lowered step by step into genuine sm_90 PTX, a custom dialect built from scratch, and the economics of who now controls the narrowest point in the AI software stack.
Three headlines, one substrate
On June 24, 2026, Qualcomm announced it would acquire Modular, the AI infrastructure company founded in 2022 by Chris Lattner and Tim Davis.
Qualcomm’s own press release gives no price; Reuters and the Wall Street Journal reported an all-stock deal worth approximately $3.9 billion, expected to close in the second half of 2026. Modular employs roughly 150 people, per the industry analysts at NAND Research.
Do the arithmetic and Qualcomm is paying something like $26 million per person for a company whose flagship language, Mojo, reached its first 1.0 beta seven weeks before the announcement, and whose compiler is not even open source yet. Qualcomm CEO Cristiano Amon justified it in one line: developers, he said, “demand a more open and modern software foundation.”
Two months earlier, in April 2026, Tesla shipped FSD v14.3. Buried in the release notes, as documented by the release trackers at Tesla Oracle, was the claim that the AI compiler and runtime had been rewritten “from the ground up with MLIR”, with Tesla attributing a 20 percent improvement in vehicle reaction time to the new stack.
And a year before that, in May 2025, NVIDIA shipped CUTLASS 4 with something unprecedented for the company: CuTe DSL, a Python interface for authoring peak-performance GPU kernels which, per NVIDIA’s own documentation, are JIT compiled through MLIR and then handed to ptxas.
A $3.9 billion acquisition by a mobile silicon giant. A safety-critical automotive inference stack. The kernel library that NVIDIA itself uses to demonstrate what Blackwell can do.
Three announcements, three companies with three entirely different agendas, one common substrate underneath. That substrate is MLIR, and if you work anywhere near GPUs, inference serving, or the economics of accelerated compute, you are running on top of it whether you know it or not.
In the previous installment of this series, How LLVM Works: The IR That Took Over Modern Computing, we told the story of the intermediate representation that unified the CPU world: one typed SSA language in the middle, every source language on one side, every instruction set on the other.
This piece is about what happened when that model met machine learning and broke, and about the infrastructure Lattner’s team built in response. MLIR is routinely described as “LLVM for AI.” It is not.
It is something stranger and more consequential: not a compiler, but a machine for building compilers, and it has quietly become the coordination point for the entire AI hardware buildout.
How to read the code in this piece
Every block of IR, TableGen, and C++ below was run through real tools before publication: mlir-opt, mlir-runner, mlir-tblgen, and GCC against the MLIR 20 headers, all from the stock LLVM 20.1.2 packages on Ubuntu 24.04.
A verified badge means the snippet parses and passes the verifier; executed means we compiled it to native code and ran it, and the output you see is what the machine printed.
The problem LLVM could not see
LLVM IR’s contract, the one that conquered the CPU world, is deliberately narrow: scalar values in SSA form, a flat control flow graph of basic blocks, loads and stores against a flat memory model, one abstraction level for everything.
That narrowness is why thirty-odd frontends and every serious instruction set could meet in the middle. It is also precisely what makes LLVM IR the wrong meeting point for machine learning.
Consider what a tensor compiler needs to know about a matrix multiplication.
That the two input buffers are dense 2-D arrays with known strides.
That the triple loop nest around it is two parallel dimensions and one reduction.
That the accumulator is transient and could live in registers or shared memory.
That the whole operation is one node in a graph whose neighbors it might profitably fuse with.
By the time a matmul has been lowered to LLVM IR, every one of those facts has been destroyed. The arrays are opaque pointers, the loop structure is a thicket of compare-and-branch, the parallelism is gone, and the fusion opportunity evaporated three layers up.
Recovering that information from LLVM IR is the auto-vectorization problem, a research area with forty years of partial credit. The lesson the ML compiler generation drew was blunt: do not recover structure. Never lose it.
The second problem was combinatorial rather than semantic. Lattner has told the origin story himself, most recently in January 2026 in “Democratizing AI Compute, Part 8,” on Modular’s blog: MLIR began in 2018 inside Google, in the middle of the TPU buildout, when the TensorFlow ecosystem had accumulated a zoo of graph representations, converters, and codegen paths that shared almost nothing.
Every framework wanted to reach every accelerator. Every accelerator vendor was rebuilding the same 80 percent of a compiler: parsers, printers, pass managers, verifiers, canonicalizers, testing harnesses.
With N frontends and M targets, the industry was signing up to build and staff something like N times M compilers.
The insight that became MLIR was to go one level more abstract than LLVM did. LLVM’s answer to fragmentation was a single fixed IR in the middle.
But ML compilation does not have one natural middle: it has a graph level, a tensor algebra level, a loop level, a vector level, a hardware intrinsic level, and programs need to descend through all of them.
So instead of shipping another fixed IR, ship the infrastructure that makes intermediate representations cheap to build, and let them all coexist in one framework, one type system, one pass manager, one textual format.
Google open sourced the project in April 2019 and donated it to the LLVM Foundation, with the code landing in the LLVM monorepo at the end of that year.
The design paper, “MLIR: Scaling Compiler Infrastructure for Domain Specific Computation” by Lattner, Mehdi Amini, Uday Bondhugula, Albert Cohen, and colleagues, appeared at CGO 2021 and is still the best single description of the architecture.
The name expands to Multi-Level Intermediate Representation, and the plural is the entire point. One artifact can hold, simultaneously and legally, a tensor-algebra view of one function, an explicit loop nest view of another, and raw hardware intrinsics for a third, with every level checked by the same verifier and transformed by the same pass infrastructure.
The rest of this article is a tour of how that actually works, at the level of real IR, because the abstractions only make sense once you have watched them move.
One idea, applied without exception
Strip away everything else and MLIR is one data structure applied with unusual discipline. The universe contains operations.
An operation has a name, takes SSA values as operands, produces values as results, carries compile-time constants called attributes, and, this is the structural leap, may contain regions: nested lists of blocks which themselves contain operations. Values have types.
That is the whole object model. There are no statements, no expressions, no special forms. Here is the smallest interesting function, in the friendly syntax you will see in every MLIR dump:
01_axpy.mlir verified: mlir-opt 20.1.2
func.func @axpy(%a: f32, %x: f32, %y: f32) -> f32 {
%0 = arith.mulf %a, %x : f32
%1 = arith.addf %0, %y : f32
return %1 : f32
}Friendly syntax is sugar. Ask mlir-opt to print the same module in generic form and the uniformity becomes visible:
mlir-opt --mlir-print-op-generic 01_axpy.mlir verified: actual tool output
"builtin.module"() ({
"func.func"() <{function_type = (f32, f32, f32) -> f32, sym_name = "axpy"}> ({
^bb0(%arg0: f32, %arg1: f32, %arg2: f32):
%0 = "arith.mulf"(%arg0, %arg1) <{fastmath = #arith.fastmath<none>}> : (f32, f32) -> f32
%1 = "arith.addf"(%0, %arg2) <{fastmath = #arith.fastmath<none>}> : (f32, f32) -> f32
"func.return"(%1) : (f32) -> ()
}) : () -> ()
}) : () -> ()Read that carefully, because it is the most important dump in this article. The module is an operation named builtin.module with one region. The function is an operation named func.func whose “body” is just a region and whose name and signature are ordinary attributes.
Even return is an operation. There are no built-in language constructs at all; the things that look built in live in a dialect that happens to be called builtin. Every structure you will meet from here on, loop nests, GPU kernels, entire schedules, is an operation containing regions containing operations, all the way down.
Regions are what LLVM IR never had. LLVM gives you one flat control flow graph per function, so a loop exists only as a cycle among basic blocks, something analyses must rediscover with dominator trees.
In MLIR, an operation like a loop contains its body as a region, so the nesting of the source program survives as nesting in the IR. A GPU kernel is an operation containing the kernel body. A module is an operation containing functions.
SSA and dominance still hold inside each region, so classic optimization theory transfers intact, but structure is now first class rather than archaeological.
Discipline this strict pays off in the machinery. Because everything is an operation, one parser, one printer, one pass manager, one multithreading strategy, and one verification framework serve every abstraction level ever built on MLIR, including ones that do not exist yet.
Extensibility is kept honest by three mechanisms.
Verifiers let each operation enforce its own structural invariants, so malformed IR dies at the boundary instead of corrupting a pass 40 steps later.
Traits declare properties like “this op has no side effects” or “operands and results share a type.”
And interfaces are the load-bearing wall: a pass written against
LoopLikeOpInterfaceorDestinationStyleOpInterfaceworks on any operation implementing the interface, including operations invented years after the pass was written by people who never met the pass author.
That is the property that makes what follows possible: generic infrastructure over an open universe of abstractions.
Dialects: a periodic table you can extend
Operations are grouped into dialects: namespaces that bundle related ops, types, and attributes together with their verifiers, canonicalization patterns, and documentation.
The prefix before the dot tells you the dialect: arith.mulf, func.func, linalg.matmul. A dialect is not a language; it is closer to a library of IR vocabulary, and dialects mix freely in one function.
The stock distribution is already a small civilization:
the in-tree roster verified: actual tool output
$ mlir-opt --show-dialects
Available Dialects (48):
acc, affine, amdgpu, amx, arith, arm_neon, arm_sme, arm_sve, async,
bufferization, builtin, cf, complex, dlti, emitc, func, gpu, index,
irdl, linalg, llvm, math, memref, mesh, ml_program, mpi, nvgpu,
nvvm, omp, pdl, pdl_interp, polynomial, ptr, quant, rocdl, scf,
shape, sparse_tensor, spirv, tensor, test, test_dyn, tosa,
transform, ub, vector, x86vector, xegpuForty-eight dialects ship in LLVM 20.1.2, and the count only grows: CPU vector extensions (amx, arm_sve, x86vector), GPU vendor exits (nvvm, rocdl, xegpu), parallel runtimes (omp, async), quantization, sparse tensors, hardware description.
Out of tree, the population explodes; we will meet the important settlers in a few sections. What matters is the shape of the traffic: programs enter at high abstraction and descend. A few landmarks, each verified:
scf, structured control flow. Loops and conditionals as region-holding operations rather than branch spaghetti. Loop-carried state is explicit SSA, threaded through iter_args:
dot product in scf verified: mlir-opt 20.1.2
// scf: structured control flow with loop-carried values
func.func @dot(%a: memref<1024xf32>, %b: memref<1024xf32>) -> f32 {
%c0 = arith.constant 0 : index
%c1 = arith.constant 1 : index
%n = arith.constant 1024 : index
%zero = arith.constant 0.0 : f32
%sum = scf.for %i = %c0 to %n step %c1
iter_args(%acc = %zero) -> (f32) {
%x = memref.load %a[%i] : memref<1024xf32>
%y = memref.load %b[%i] : memref<1024xf32>
%m = arith.mulf %x, %y : f32
%s = arith.addf %acc, %m : f32
scf.yield %s : f32
}
return %sum : f32
}affine, the polyhedral inheritance. Same loops, harsher rules: bounds and subscripts must be affine expressions of the induction variables. In exchange, dependence analysis becomes decidable.
The compiler can prove that iterations are independent, that tiling is legal, that this access pattern never aliases that one, instead of guessing:
1-D stencil in affine verified: mlir-opt 20.1.2
// affine: loops + accesses restricted to affine expressions, so the
// compiler can *prove* dependence facts instead of guessing
func.func @blur1d(%in: memref<1026xf32>, %out: memref<1024xf32>) {
affine.for %i = 0 to 1024 {
%l = affine.load %in[%i] : memref<1026xf32>
%c = affine.load %in[%i + 1] : memref<1026xf32>
%r = affine.load %in[%i + 2] : memref<1026xf32>
%s0 = arith.addf %l, %c : f32
%s1 = arith.addf %s0, %r : f32
affine.store %s1, %out[%i] : memref<1024xf32>
}
return
}tensor and memref, the two memories. This pair encodes the single most useful distinction in the whole stack.
A tensor is an immutable SSA value: no address, no aliasing, safe to reorder and fuse aggressively, the natural currency of ML graphs.
A memref is a view of actual memory: base pointer, sizes, strides, and, critically for this audience, an address space. The types are expressive enough to carry real kernel engineering:
types that will look familiar verified: mlir-opt 20.1.2
func.func private @smem_tile(%t: memref<64x64xf16, strided<[68, 1]>, 3>)
func.func private @kv_page(%p: tensor<16x8x128xbf16>)The first is a 64x64 half-precision tile with a padded row stride of 68 elements, the classic plus-four skew that kills shared memory bank conflicts, living in address space 3: GPU shared memory.
The second is a bf16 paged-KV block of the kind every serving engine slings around. MLIR’s type system says these things natively; nothing here is a comment or a convention.
vector, the portable SIMD layer. Fixed-shape virtual registers and the operations that move them: vector.fma, transfers, contractions, shuffles. Target independent until the last moment, then peeled into NEON, AVX-512, or GPU instructions:
an 8-wide fused multiply-add verified: mlir-opt 20.1.2
// vector: SIMD/SIMT-width-explicit compute on virtual registers
func.func @vec_fma(%a: vector<8xf32>, %b: vector<8xf32>, %c: vector<8xf32>) -> vector<8xf32> {
%0 = vector.fma %a, %b, %c : vector<8xf32>
return %0 : vector<8xf32>
}And beneath everything, the exits: the gpu dialect for launch grids and kernels, nvvm and rocdl wrapping the vendors’ intrinsics, and the llvm dialect, which models LLVM IR itself as just another MLIR dialect so that the handoff to the old empire is an ordinary rewrite rather than a file format negotiation.
One table to orient the descent:
Read the right-hand column top to bottom and you have the core operating principle of the whole system, the one we will watch in motion next: information is only destroyed downward, so every optimization should run at the highest level where its enabling facts still exist.
Fusion happens in tensor algebra, where aliasing cannot exist. Tiling happens where loops are structured objects. Bank conflict avoidance happens where address spaces are types.
By the time you reach LLVM, you are not optimizing anymore; you are translating.
linalg: the algebra at the waist
If MLIR has a center of gravity, it is the linalg dialect: the tensor-algebra layer where most serious codegen pipelines, IREE’s, Modular’s, the upstream CPU flows, do their thinking.
Its design principle is the anti-autovectorizer stance taken to its conclusion. Instead of encoding a matmul as loops and hoping to rediscover its nature, linalg encodes the nature and treats loops as one of several possible spellings, generated on demand.
The workhorse is linalg.generic, which describes a perfectly nested computation with three pieces of data. First, indexing maps: one affine map per operand saying which element each point in the iteration space touches.
Second, iterator types: each loop dimension is declared parallel or reduction, which is the fact autovectorizers spend their lives failing to prove. Third, a region holding the scalar payload, the body computed at every point.
Everything the dialect knows, it knows structurally. A matmul’s maps are (m, n, k) -> (m, k), (m, n, k) -> (k, n), (m, n, k) -> (m, n) over iterators [parallel, parallel, reduction]; named ops like linalg.matmul are simply memorable spellings of specific generics.
Here is the payload we will spend the rest of the article compiling, a real inference fragment: a linear layer with fused bias and ReLU.
02_mlp.mlir: matmul + fused bias/ReLU on tensors verified: mlir-opt 20.1.2
#id2d = affine_map<(m, n) -> (m, n)>
#bcast = affine_map<(m, n) -> (n)>
func.func @mlp_block(%x: tensor<64x512xf32>,
%w: tensor<512x256xf32>,
%b: tensor<256xf32>) -> tensor<64x256xf32> {
%c0 = arith.constant 0.0 : f32
// Materialize the destination and zero-init the accumulator.
%empty = tensor.empty() : tensor<64x256xf32>
%acc = linalg.fill ins(%c0 : f32)
outs(%empty : tensor<64x256xf32>) -> tensor<64x256xf32>
// y = x @ w
%mm = linalg.matmul ins(%x, %w : tensor<64x512xf32>, tensor<512x256xf32>)
outs(%acc : tensor<64x256xf32>) -> tensor<64x256xf32>
// out = max(y + b, 0), one fused elementwise op
%out = linalg.generic
{indexing_maps = [#id2d, #bcast, #id2d],
iterator_types = ["parallel", "parallel"]}
ins(%mm, %b : tensor<64x256xf32>, tensor<256xf32>)
outs(%empty : tensor<64x256xf32>) {
^bb0(%y: f32, %bias: f32, %o: f32):
%sum = arith.addf %y, %bias : f32
%relu = arith.maximumf %sum, %c0 : f32
linalg.yield %relu : f32
} -> tensor<64x256xf32>
return %out : tensor<64x256xf32>
}Three details deserve attention. The bias’s indexing map, (m, n) -> (n), expresses broadcasting as pure index algebra: no replicated buffer, no runtime rule, just a map that ignores m.
The elementwise op fuses the add and the clamp into one region, because a region can hold any scalar computation, which is how epilogue fusion is spelled at this level.
And every linalg op writes into an explicit outs operand it also returns: destination-passing style. On immutable tensors the destination looks redundant, and semantically it is.
It exists as a standing declaration of where results could live, which is exactly the hint that lets bufferization later plan in-place updates instead of solving a global aliasing puzzle. Keep an eye on that %empty; it is about to earn its keep.
The walkthrough: one MLP block to sm_90 PTX
Now the demonstration the whole article is built around. We will take that block and walk it down the ladder one verified step at a time, watching what each level adds and what it forgets.
This is the same journey your PyTorch graphs take through Triton, that JAX programs take through XLA, that Mojo kernels take through MAX; we are just taking it on foot, with the intermediate files open.
Step 1. Tiling, written as IR
The first performance decision on any matrix multiply is blocking it for the memory hierarchy. In most compilers, tile sizes hide inside a C++ heuristic.
In MLIR, the transformation itself can be scripted in IR, using the transform dialect: a schedule language whose operations manipulate other operations. We append this to the file:
03_tile.mlir, the schedule part verified: mlir-opt 20.1.2
// ---- Schedule: also IR, in the transform dialect ----
module attributes {transform.with_named_sequence} {
transform.named_sequence @__transform_main(%root: !transform.any_op {transform.readonly}) {
%mm = transform.structured.match ops{["linalg.matmul"]} in %root
: (!transform.any_op) -> !transform.any_op
%tiled, %loops:3 = transform.structured.tile_using_for %mm tile_sizes [16, 32, 64]
: (!transform.any_op) -> (!transform.any_op, !transform.any_op, !transform.any_op, !transform.any_op)
transform.yield
}
}The handles typed !transform.any_op are SSA values whose runtime contents are payload operations: the schedule matches every linalg.matmul, then tiles each one 16 by 32 by 64.
Run mlir-opt --transform-interpreter and the payload function comes back rewritten:
output: the tiled matmul verified: actual tool output, constants elided as marked
func.func @matmul(%arg0: tensor<64x512xf32>, %arg1: tensor<512x256xf32>, %arg2: tensor<64x256xf32>) -> tensor<64x256xf32> {
// ... index constants (0, 16, 32, 64, 256, 512) elided ...
%0 = scf.for %arg3 = %c0 to %c64 step %c16 iter_args(%arg4 = %arg2) -> (tensor<64x256xf32>) {
%1 = scf.for %arg5 = %c0_0 to %c256 step %c32 iter_args(%arg6 = %arg4) -> (tensor<64x256xf32>) {
%2 = scf.for %arg7 = %c0_1 to %c512 step %c64_2 iter_args(%arg8 = %arg6) -> (tensor<64x256xf32>) {
%extracted_slice = tensor.extract_slice %arg0[%arg3, %arg7] [16, 64] [1, 1] : tensor<64x512xf32> to tensor<16x64xf32>
%extracted_slice_3 = tensor.extract_slice %arg1[%arg7, %arg5] [64, 32] [1, 1] : tensor<512x256xf32> to tensor<64x32xf32>
%extracted_slice_4 = tensor.extract_slice %arg8[%arg3, %arg5] [16, 32] [1, 1] : tensor<64x256xf32> to tensor<16x32xf32>
%3 = linalg.matmul ins(%extracted_slice, %extracted_slice_3 : tensor<16x64xf32>, tensor<64x32xf32>) outs(%extracted_slice_4 : tensor<16x32xf32>) -> tensor<16x32xf32>
%inserted_slice = tensor.insert_slice %3 into %arg8[%arg3, %arg5] [16, 32] [1, 1] : tensor<16x32xf32> into tensor<64x256xf32>
scf.yield %inserted_slice : tensor<64x256xf32>
}
scf.yield %2 : tensor<64x256xf32>
}
scf.yield %1 : tensor<64x256xf32>
}
return %0 : tensor<64x256xf32>
}One schedule op became a three-deep scf.for nest. tensor.extract_slice carves 16x64 and 64x32 tiles from the operands, an inner linalg.matmul, still the full structured op, now shaped 16x64 times 64x32, computes each 16x32 block, and tensor.insert_slice threads the result through the loop-carried tensor.
Notice what did not happen: we have loops, but we did not lose the algebra. The inner op is still a matmul with its maps and iterator types intact, still available for vectorization, for lowering to tensor core intrinsics, or for another round of tiling.
Structure survives the transformation because the transformation is structure-aware.
Step 2. Fusion: FlashAttention thinking, six lines long
Tiling one op is table stakes. The move that defines modern kernel engineering, compute the producer inside the consumer’s tile so intermediates never round-trip through HBM, is the same idea that makes FlashAttention FlashAttention.
Here it is as a schedule: tile the bias/ReLU consumer into a parallel grid, then pull the matmul into each tile.
04_fuse.mlir, the schedule part verified: mlir-opt 20.1.2
module attributes {transform.with_named_sequence} {
transform.named_sequence @__transform_main(%root: !transform.any_op {transform.readonly}) {
// 1. Tile the *consumer* (bias+relu) into a parallel grid of 16x32 tiles.
%gen = transform.structured.match ops{["linalg.generic"]} in %root
: (!transform.any_op) -> !transform.any_op
%tiled, %grid = transform.structured.tile_using_forall %gen tile_sizes [16, 32]
: (!transform.any_op) -> (!transform.any_op, !transform.any_op)
// 2. Pull the matmul producer inside each tile of that grid.
%mm = transform.structured.match ops{["linalg.matmul"]} in %root
: (!transform.any_op) -> !transform.any_op
%fused, %grid2 = transform.structured.fuse_into_containing_op %mm into %grid
: (!transform.any_op, !transform.any_op) -> (!transform.any_op, !transform.any_op)
transform.yield
}
}output after --transform-interpreter -canonicalize -cse verified: actual tool output
#map = affine_map<(d0) -> (d0 * 16)>
#map1 = affine_map<(d0) -> (d0 * 32)>
#map2 = affine_map<(d0, d1) -> (d0, d1)>
#map3 = affine_map<(d0, d1) -> (d1)>
func.func @mlp_block(%arg0: tensor<64x512xf32>, %arg1: tensor<512x256xf32>, %arg2: tensor<256xf32>) -> tensor<64x256xf32> {
%cst = arith.constant 0.000000e+00 : f32
%0 = tensor.empty() : tensor<64x256xf32>
%1 = linalg.fill ins(%cst : f32) outs(%0 : tensor<64x256xf32>) -> tensor<64x256xf32>
%2 = scf.forall (%arg3, %arg4) in (4, 8) shared_outs(%arg5 = %0) -> (tensor<64x256xf32>) {
%3 = affine.apply #map(%arg3)
%4 = affine.apply #map1(%arg4)
%extracted_slice = tensor.extract_slice %arg0[%3, 0] [16, 512] [1, 1] : tensor<64x512xf32> to tensor<16x512xf32>
%extracted_slice_0 = tensor.extract_slice %arg1[0, %4] [512, 32] [1, 1] : tensor<512x256xf32> to tensor<512x32xf32>
%extracted_slice_1 = tensor.extract_slice %1[%3, %4] [16, 32] [1, 1] : tensor<64x256xf32> to tensor<16x32xf32>
%5 = linalg.matmul ins(%extracted_slice, %extracted_slice_0 : tensor<16x512xf32>, tensor<512x32xf32>) outs(%extracted_slice_1 : tensor<16x32xf32>) -> tensor<16x32xf32>
%extracted_slice_2 = tensor.extract_slice %arg2[%4] [32] [1] : tensor<256xf32> to tensor<32xf32>
%extracted_slice_3 = tensor.extract_slice %arg5[%3, %4] [16, 32] [1, 1] : tensor<64x256xf32> to tensor<16x32xf32>
%6 = linalg.generic {indexing_maps = [#map2, #map3, #map2], iterator_types = ["parallel", "parallel"]} ins(%5, %extracted_slice_2 : tensor<16x32xf32>, tensor<32xf32>) outs(%extracted_slice_3 : tensor<16x32xf32>) {
^bb0(%in: f32, %in_4: f32, %out: f32):
%7 = arith.addf %in, %in_4 : f32
%8 = arith.maximumf %7, %cst : f32
linalg.yield %8 : f32
} -> tensor<16x32xf32>
scf.forall.in_parallel {
tensor.parallel_insert_slice %6 into %arg5[%3, %4] [16, 32] [1, 1] : tensor<16x32xf32> into tensor<64x256xf32>
}
}
return %2 : tensor<64x256xf32>
}The consumer became an scf.forall, a parallel-by-construction grid of 4 by 8 tiles, and inside each tile sits a private 16x512 by 512x32 matmul feeding the fused bias-and-clamp directly.
The 64x256 intermediate matrix that used to exist between the two ops is gone from the program: each tile’s slice is produced, consumed, and discarded in place. canonicalize and cse then deleted the original full-size matmul as dead code.
This is a working, verified skeleton of producer fusion, the transformation entire serving stacks are built around, expressed in six lines of schedule that a performance engineer can read, diff, and review like any other code.
Step 3. Bufferization: the value world ends here
Everything so far lived on immutable tensors, where fusion and reordering are trivially safe because aliasing cannot be expressed. Hardware, unfortunately, sells mutable bytes.
The crossing is one-shot bufferization, a whole-module analysis that assigns every tensor a buffer while proving where it can reuse memory instead of copying. Run it on the unfused MLP and look at what comes back:
mlir-opt -one-shot-bufferize=”bufferize-function-boundaries” verified: actual tool output
#map = affine_map<(d0, d1) -> (d0, d1)>
#map1 = affine_map<(d0, d1) -> (d1)>
module {
func.func @mlp_block(%arg0: memref<64x512xf32, strided<[?, ?], offset: ?>>, %arg1: memref<512x256xf32, strided<[?, ?], offset: ?>>, %arg2: memref<256xf32, strided<[?], offset: ?>>) -> memref<64x256xf32> {
%cst = arith.constant 0.000000e+00 : f32
%alloc = memref.alloc() {alignment = 64 : i64} : memref<64x256xf32>
linalg.fill ins(%cst : f32) outs(%alloc : memref<64x256xf32>)
linalg.matmul ins(%arg0, %arg1 : memref<64x512xf32, strided<[?, ?], offset: ?>>, memref<512x256xf32, strided<[?, ?], offset: ?>>) outs(%alloc : memref<64x256xf32>)
linalg.generic {indexing_maps = [#map, #map1, #map], iterator_types = ["parallel", "parallel"]} ins(%alloc, %arg2 : memref<64x256xf32>, memref<256xf32, strided<[?], offset: ?>>) outs(%alloc : memref<64x256xf32>) {
^bb0(%in: f32, %in_0: f32, %out: f32):
%0 = arith.addf %in, %in_0 : f32
%1 = arith.maximumf %0, %cst : f32
linalg.yield %1 : f32
}
%cast = memref.cast %alloc : memref<64x256xf32> to memref<64x256xf32, strided<[?, ?], offset: ?>>
return %alloc : memref<64x256xf32>
// ...Tensors became memrefs; function boundaries acquired conservative strided layouts so callers can pass any compatible view.
But the line to study is the fused elementwise: it now reads from %alloc and writes into %alloc. The analysis proved the bias/ReLU can execute in place on the matmul’s buffer, so the second 64KB allocation that a naive translation would emit simply does not exist.
That is destination-passing style paying out: because every linalg op declared its destination back in value land, the planner had the aliasing story handed to it.
This pass, more than any other, is where MLIR’s two-world design earns its complexity: algebra above the line, memory below it, and one well-tested analysis holding the border.
Step 4. Proof of life: the IR actually runs
Before descending to GPU intrinsics, a checkpoint that separates this article from a whiteboard exercise.
The remaining distance to a CPU binary is mechanical: linalg to loops, loops to branches, memrefs and arithmetic to the llvm dialect, then JIT.
We wrapped the MLP in a tiny main with known inputs, chose a bias of [-10, 1] so the ReLU visibly clamps, and ran the pipeline:
the full descent, one command executed: mlir-runner, LLVM 20.1.2
$ mlir-opt mlp_run.mlir \
-one-shot-bufferize="bufferize-function-boundaries" \
-convert-linalg-to-loops \
-expand-strided-metadata -lower-affine \
-convert-scf-to-cf -finalize-memref-to-llvm \
-convert-arith-to-llvm -convert-cf-to-llvm -convert-func-to-llvm \
-reconcile-unrealized-casts \
| mlir-runner -e main -entry-point-result=void \
-shared-libs=$LLVM/lib/libmlir_runner_utils.so.20.1,$LLVM/lib/libmlir_c_runner_utils.so.20.1what the machine printed
Unranked Memref base@ = 0x55b27c52f500 rank = 2 offset = 0 sizes = [2, 2] strides = [2, 1] data =
[[0, 6],
[0, 12]]By hand: x@w is [[4, 5], [10, 11]]; add the bias to get [[-6, 6], [0, 12]]; clamp at zero for [[0, 6], [0, 12]]. The machine agrees.
Every abstraction in this article compiles to electrons, and the fused ReLU is provably doing its job on the negative entries.
Step 5. The GPU descent, ending in genuine PTX
Now the branch this audience came for. Same matmul, memref form, aimed at a GPU.
First, expose the parallelism that linalg has known about all along: -convert-linalg-to-parallel-loops turns the two parallel iterators into an scf.parallel over (m, n) with the reduction as an ordinary inner loop:
parallel loops verified: actual tool output
func.func @matmul(%arg0: memref<64x512xf32>, %arg1: memref<512x256xf32>, %arg2: memref<64x256xf32>) {
%c0 = arith.constant 0 : index
%c64 = arith.constant 64 : index
%c1 = arith.constant 1 : index
%c256 = arith.constant 256 : index
%c512 = arith.constant 512 : index
scf.parallel (%arg3, %arg4) = (%c0, %c0) to (%c64, %c256) step (%c1, %c1) {
scf.for %arg5 = %c0 to %c512 step %c1 {
%0 = memref.load %arg0[%arg3, %arg5] : memref<64x512xf32>
%1 = memref.load %arg1[%arg5, %arg4] : memref<512x256xf32>
%2 = memref.load %arg2[%arg3, %arg4] : memref<64x256xf32>
%3 = arith.mulf %0, %1 : f32
%4 = arith.addf %2, %3 : f32
memref.store %4, %arg2[%arg3, %arg4] : memref<64x256xf32>
}
scf.reduce
}
return
}
}Next, the mapping decision: -gpu-map-parallel-loops -convert-parallel-loops-to-gpu assigns parallel dimensions to the grid and materializes a launch:
a kernel is born verified: actual tool output
gpu.launch blocks(%arg3, %arg4, %arg5) in (%arg9 = %0, %arg10 = %1, %arg11 = %c1_0) threads(%arg6, %arg7, %arg8) in (%arg12 = %c1_0, %arg13 = %c1_0, %arg14 = %c1_0) {
%2 = affine.apply #map1(%arg3)[%c1, %c0]
%3 = affine.apply #map1(%arg4)[%c1, %c0]
scf.for %arg15 = %c0 to %c512 step %c1 {
%4 = memref.load %arg0[%2, %arg15] : memref<64x512xf32>
%5 = memref.load %arg1[%arg15, %3] : memref<512x256xf32>
%6 = memref.load %arg2[%2, %3] : memref<64x256xf32>
%7 = arith.mulf %4, %5 : f32
%8 = arith.addf %6, %7 : f32
memref.store %8, %arg2[%2, %3] : memref<64x256xf32>
}
// ... gpu.terminatorBe honest about what we are looking at: a 64 by 256 grid of blocks with a single thread each, the most naive mapping a GPU has ever been insulted with.
Production pipelines tile before mapping, so blocks get tiles and threads get elements, then layer on shared memory staging and tensor core ops from the nvgpu dialect (mma.sync and friends).
We are deliberately taking the unoptimized path because the plumbing, not the schedule, is today’s subject; later sections covers who builds the good schedules. -gpu-kernel-outlining then splits host from device, the step every CUDA programmer performs mentally when writing __global__:
outlined device module verified: actual tool output, trimmed as marked
gpu.module @matmul_kernel {
gpu.func @matmul_kernel(%arg0: index, %arg1: index, %arg2: memref<64x512xf32>, %arg3: memref<512x256xf32>, %arg4: memref<64x256xf32>, %arg5: index) kernel attributes {known_block_size = array<i32: 1, 1, 1>} {
%block_id_x = gpu.block_id x
%block_id_y = gpu.block_id y
%thread_id_x = gpu.thread_id x
%thread_id_y = gpu.thread_id y
%0 = affine.apply #map1(%block_id_x)[%arg0, %arg1]
%1 = affine.apply #map1(%block_id_y)[%arg0, %arg1]
// ... same loop body, now reading gpu.block_id instead of loop IVsThe loop body is unchanged, but the block indices now arrive from gpu.block_id instead of loop induction variables, and the kernel records its launch invariants (known_block_size) as attributes for later passes to exploit.
From here, -convert-gpu-to-nvvm rewrites device code into the llvm dialect sprinkled with NVIDIA’s intrinsics:
nvvm: the vendor boundary verified: excerpt of actual tool output
gpu.module @matmul_kernel {
// signature abbreviated: 24 pointer and index arguments
llvm.func @matmul_kernel(%arg0: i64, ..., %arg23: i64)
attributes {gpu.kernel, gpu.known_block_size = array<i32: 1, 1, 1>,
nvvm.kernel, nvvm.maxntid = array<i32: 1, 1, 1>} {
// ...
%24 = nvvm.read.ptx.sreg.ctaid.x : i32
%25 = llvm.sext %24 : i32 to i64
%26 = nvvm.read.ptx.sreg.ctaid.y : i32
// ... address arithmetic in llvm dialect: getelementptr, mul, add ...
}
}gpu.block_id x became nvvm.read.ptx.sreg.ctaid.x, a name any CUDA disassembly veteran will greet like an old acquaintance: the special register file, now visible in the IR.
Finally, the packaged pipeline --gpu-lower-to-nvvm-pipeline="cubin-chip=sm_90 cubin-format=isa" runs the whole descent and invokes LLVM’s NVPTX backend, embedding the result in a gpu.binary op. Extract it and you are holding 90 lines of the real thing:
out/matmul.ptx generated: LLVM 20.1.2 NVPTX backend, sm_90
//
// Generated by LLVM NVPTX Back-End
//
.version 7.8
.target sm_90
.address_size 64
// .globl matmul_kernel
// ... prologue: bounds check, address setup ...
$L__BB0_2:
ld.global.f32 %f4, [%rd35];
ld.global.f32 %f5, [%rd34];
mul.rn.f32 %f6, %f4, %f5;
add.rn.f32 %f7, %f7, %f6;
st.global.f32 [%rd6], %f7;
add.s64 %rd36, %rd36, %rd17;
add.s64 %rd35, %rd35, %rd8;
add.s64 %rd34, %rd34, %rd10;
setp.lt.s64 %p2, %rd36, %rd19;
@%p2 bra $L__BB0_2;
// ... epilogue ... (90 lines total)There is a final teaching moment hiding in that inner loop, and it is the thesis of this article restated by the machine. Look at st.global.f32: the accumulator is written back to global memory every iteration of k.
The value lives happily in register %f7, yet the store stays, because at this level nothing can prove the output buffer does not alias the inputs, so the backend must assume a reader might observe every intermediate sum.
Four levels up, in linalg on tensors, that fact was free: tensors cannot alias. We destroyed the information on the way down, exactly as the ladder predicted, and the machine code is paying for it, one redundant HBM transaction per multiply.
Every dollar of kernel engineering ever spent is, in some form, the cost of putting facts back that a lowering threw away. MLIR’s entire bet is that it is cheaper to never throw them away.
Be precise about the distance between this demonstration and a kernel you would actually deploy, because the gap is the whole subject of the next section.
A production pipeline tiles before it maps to hardware, so a block owns a tile and a thread owns a register-sized sliver rather than the one-thread-per-block insult we emitted.
It stages the operand tiles through the padded shared-memory memref from section 4, turning the plus-four skew into live bank-conflict avoidance.
And instead of scalar mul.rn.f32 plus that redundant store, it lowers the inner tile op through the nvgpu dialect into mma.sync or, on Hopper, wgmma, so the loop body becomes a warp-level matrix instruction feeding tensor cores, with the accumulator held in registers across the entire reduction and written back exactly once.
Here is the point that matters for understanding the whole system: every one of those improvements is a different schedule over the identical set of passes we just ran. The naive descent and the peak one share their entire lowering machinery; what separates 2 percent of peak from 90 percent is the sequence of tiling, fusion, and mapping decisions applied on the way down.
Which is why the ability to write that sequence as a first-class, reviewable, searchable artifact, rather than bury it in compiler C++, is not a curiosity. It is the ballgame.
The schedule is also IR
Steps 1 and 2 of the walkthrough smuggled in the most radical idea in modern compiler engineering, so let us stop and look at it directly.
Classically, a compiler’s transformations are opaque C++ controlled by flags and heuristics; if the heuristic tiles your attention kernel wrong, your options are a rebuild or a prayer.
Halide’s great insight, back in 2012, was to split the algorithm from the schedule and make the schedule a first-class program. MLIR’s transform dialect imports that split into the compiler infrastructure itself: the schedule is IR, in the same file format, checked by the same verifier, with SSA handles pointing at the payload operations it manipulates.
The consequences compound. A schedule that is data can be diffed and code reviewed: the six lines in step 2 are a reviewable artifact in a way a C++ pass never is.
It can be searched: an autotuner can emit thousands of candidate schedules, run the interpreter, and measure, without recompiling the compiler; the transform interpreter is exactly the substrate you would build a kernel search system on.
It can be shipped per workload: IREE has used transform-dialect scripts to pin codegen strategies for specific dispatches, which is how you get reproducible kernels out of a general compiler.
And it fails loudly: apply a tiling to an op that cannot be tiled and the interpreter reports a definite error against a definite handle, rather than silently generating slow code.
For a CUDA-native audience the honest framing is this: the transform dialect is the compiler conceding that you, the performance engineer, sometimes know better, and giving you a typed, verifiable console instead of a fork of the codebase.
The upstream vocabulary already covers the daily verbs, tile_using_for, tile_using_forall, fuse_into_containing_op, vectorize, pad, plus pattern application and PDL-based matching for everything else.
It is not the default path for most users and may never be; it is the expert path, and its existence changes what the expert path costs.
Build a dialect before lunch
The claim that MLIR makes IRs cheap deserves a demonstration with a bill attached. Suppose we run inference infrastructure and want the compiler to reason about our world: quantized weights, dequantization groups, fused epilogues.
In a classical compiler, adding a first-class operation means touching the parser, the printer, the verifier, the serializer, and a dozen switch statements.
In MLIR it means describing the op once, declaratively, in TableGen’s Operation Definition Specification, and generating the rest. Here is a real op for a real concern, weight-only quantized matmul, in 28 lines:
serve_ops.td verified: mlir-tblgen 20.1.2
include "mlir/IR/OpBase.td"
include "mlir/Interfaces/SideEffectInterfaces.td"
def Serve_Dialect : Dialect {
let name = "serve";
let summary = "Ops for LLM serving kernels";
let cppNamespace = "::mlir::serve";
}
class Serve_Op<string mnemonic, list<Trait> traits = []>
: Op<Serve_Dialect, mnemonic, traits>;
def Serve_DequantMatmulOp : Serve_Op<"dequant_matmul", [Pure]> {
let summary = "matmul with fused groupwise weight dequantization";
let description = [{
Computes out = act @ dequant(weights, scales) without ever
materializing the dequantized weight tensor in memory.
}];
let arguments = (ins
TensorOf<[F16]>:$act,
TensorOf<[I8]>:$weights,
TensorOf<[F16]>:$scales,
I64Attr:$group_size
);
let results = (outs TensorOf<[F16]>:$out);
let assemblyFormat =
"$act `,` $weights `,` $scales attr-dict `:` functional-type(operands, results)";
}Feed it to mlir-tblgen and the machinery materializes:
the leverage, measured verified: actual tool output
$ mlir-tblgen -gen-op-decls -I /usr/lib/llvm-20/include serve_ops.td > ServeOps.h.inc
$ mlir-tblgen -gen-op-defs -I /usr/lib/llvm-20/include serve_ops.td > ServeOps.cpp.inc
# 28 lines of ODS in, 259 lines of generated C++ declarations out:
class DequantMatmulOp;
...Twenty-eight declarative lines became 259 lines of generated C++ declarations, plus definitions: typed accessors (getAct(), getGroupSize()), builders, parser and printer honoring our assemblyFormat, and verifier scaffolding enforcing the type constraints, an f16 activation tensor, i8 weights, f16 scales, or the IR does not construct.
Nine-to-one leverage on the boilerplate, and every generated line is the same battle-tested code path the 48 upstream dialects use.
This is what “infrastructure” means concretely: our bespoke serving op gets round-trippable syntax, verification, and pass-manager citizenship for the price of a lunch break, which is why every hardware startup in section 9 could afford a real compiler at all.
Ops are half a dialect; the other half is transformations, and those are written in C++ against the pattern infrastructure. Here is a complete, compiling pass with a pattern this audience will feel in their DRAM bills: detect a matmul whose accumulator is a fresh zero fill, and tag it so a later lowering can emit a beta-equals-zero GEMM instead of a memset followed by a GEMM, saving one full write-and-read pass over the output through HBM:
TagZeroInit.cpp verified: g++ -fsyntax-only against MLIR 20 headers, exit 0
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/Linalg/IR/Linalg.h"
#include "mlir/IR/BuiltinOps.h"
#include "mlir/IR/PatternMatch.h"
#include "mlir/Pass/Pass.h"
#include "mlir/Transforms/GreedyPatternRewriteDriver.h"
using namespace mlir;
namespace {
// Match linalg.fill(+0.0) feeding a linalg.matmul accumulator and tag
// the matmul, so a later lowering can emit a beta = 0 GEMM instead of
// a memset followed by a GEMM (one less trip through HBM).
struct TagZeroInitMatmul : OpRewritePattern<linalg::MatmulOp> {
using OpRewritePattern::OpRewritePattern;
LogicalResult matchAndRewrite(linalg::MatmulOp op,
PatternRewriter &rewriter) const override {
if (op->hasAttr("serve.zero_init"))
return failure(); // already rewritten: stop, or loop forever
auto fill =
op.getDpsInitOperand(0)->get().getDefiningOp<linalg::FillOp>();
if (!fill)
return failure();
auto cst =
fill.getInputs()[0].getDefiningOp<arith::ConstantFloatOp>();
if (!cst || !cst.value().isPosZero())
return failure();
rewriter.modifyOpInPlace(op, [&] {
op->setAttr("serve.zero_init", rewriter.getUnitAttr());
});
return success();
}
};
struct TagZeroInitPass
: PassWrapper<TagZeroInitPass, OperationPass<ModuleOp>> {
MLIR_DEFINE_EXPLICIT_INTERNAL_INLINE_TYPE_ID(TagZeroInitPass)
StringRef getArgument() const final { return "serve-tag-zero-init"; }
StringRef getDescription() const final {
return "Tag matmuls whose accumulator is a fresh zero fill";
}
void runOnOperation() override {
RewritePatternSet patterns(&getContext());
patterns.add<TagZeroInitMatmul>(&getContext());
if (failed(applyPatternsGreedily(getOperation(), std::move(patterns))))
signalPassFailure();
}
};
} // namespaceThe anatomy generalizes to every pattern you will ever write. Match structurally by walking use-def edges (getDefiningOp through the destination operand, courtesy of destination-passing style).
Guard exhaustively, including against your own previous rewrite, because the greedy driver reapplies patterns to fixpoint and an unguarded pattern is an infinite loop.
Mutate only through the rewriter, never behind its back, so the driver’s worklist stays coherent. Fifty-ish lines, and it composes with every canonicalization and lowering upstream ships.
The last pillar of the workflow is cultural: everything above gets a FileCheck test, a .mlir file whose RUN line invokes the pass and whose CHECK lines assert on the output IR.
Because IR is text with stable syntax, transformation tests are diffs, reviewable by humans and replayable by CI. Teams that build on MLIR inherit this discipline for free, and it shows: it is the reason a project like
Triton can refactor its entire layout system without regressing a thousand kernels silently.
The honest total cost of a production dialect is of course higher than lunch, the semantics and the verifier discipline are the real work, but the floor has moved by an order of magnitude, and floors are what determine who gets to play.
Where MLIR is hiding in your stack
Infrastructure succeeds when it disappears. Seven years after open sourcing, MLIR has disappeared into more of the AI stack than almost anyone tracks, so let us make the map explicit.
Triton, which means PyTorch. When OpenAI rebuilt Triton on MLIR in 2022, the language stopped being a research artifact and became a compiler stack.
Python is traced into the triton dialect (TTIR), a target-independent tile program; conversion to the triton_gpu dialect (TTGIR) then makes the performance-critical decision explicit by attaching a layout encoding to every tensor, describing exactly which thread of which warp owns which element, with the LinearLayout framework giving those mappings a uniform algebra; vendor sub-dialects handle Hopper and CDNA specifics before the drop into the llvm dialect and PTX or AMDGCN.
Since PyTorch 2 made Inductor the default compile path and Inductor emits Triton for GPU graphs, the practical consequence is that most compiled PyTorch GPU workloads on the planet route through MLIR today, whether their owners have heard of it or not.
And in 2025 the Triton team surfaced the lower rungs of its own ladder as a product: Gluon, documented in the Triton repository, hands expert users the TTGIR-level controls, layouts, shared memory, warp specialization, while reusing the same dialect stack underneath.
Which means your serving engine, too. The inference stacks this publication’s readers actually operate, vLLM and SGLang chief among them, are not usually thought of as MLIR consumers, but trace the dependency and the connection is direct.
Both lean heavily on Triton for their fused and custom kernels: fused RMSNorm and RoPE, fused MoE routing, quantized GEMM epilogues, and increasingly the attention paths through backends in the FlashInfer lineage that mix hand-written CUDA with Triton-generated variants.
Every one of those Triton kernels is a program descending through TTIR and TTGIR, which is to say through MLIR, before it reaches ptxas. The abstraction that a serving engine calls “a Triton kernel” is, one layer down, exactly the layout-annotated dialect IR.
And the data structure those engines sling most obsessively, the paged KV cache, is precisely the strided, address-spaced memref from section 4 made concrete: a block table indexing fixed-size buffers, the same view-of-memory abstraction the type system was built to name.
When you profile a vLLM deployment on a B300 and find time in a fused decode kernel, you are, whether the dashboard says so or not, looking at the output of the compiler stack this article has been dissecting. The waist is not upstream of your infrastructure; it is inside it.
NVIDIA, voluntarily. CUTLASS 4’s CuTe DSL lets engineers author kernels in Python that, per NVIDIA’s documentation, JIT through MLIR into optimized code finished by ptxas, with the CuTe layout algebra as the programming model.
Read that strategically: the company with the most to gain from CUDA C++ lock-in decided its flagship kernel library’s productivity layer should be built on the industry’s shared compiler substrate.
NVIDIA is not ceding the moat, ptxas and SASS remain closed, but it has conceded where the moat is not: the language layer above the driver is now contested ground, and NVIDIA chose to contest it with MLIR rather than against it.
Google, the origin. XLA remains the TPU’s compiler, and since 2022 its portable front door is StableHLO, an MLIR dialect with versioned serialization that lets JAX, TensorFlow, and PyTorch programs travel between compilers without marrying one.
IREE, under the OpenXLA umbrella, is the fully MLIR-native end-to-end: StableHLO or Torch in, linalg in the middle, compiled artifacts for CPU, CUDA, ROCm, Vulkan, and mobile out. It is the closest thing to the textbook hourglass shipped as a product, and its codegen is where much of the transform-dialect and structured-codegen machinery in this article was hardened.
The silicon insurgents, which is the actual story. Look at who builds their entire software identity on this infrastructure. Tenstorrent’s tt-mlir stack, public on GitHub since 2024, compiles graphs through TTIR into its TTNN library ops for Wormhole and Blackhole silicon.
AMD ships MLIR-AIE, its 1.2 release landing in January 2026 per Phoronix, as the low-level programming layer for the NPUs in every Ryzen AI laptop. Qualcomm’s compiler group published Hexagon-MLIR on arXiv this spring: a stack that ingests, notably, both PyTorch graphs and Triton kernels and lowers them onto Hexagon NPUs.
Tesla, per the FSD v14.3 release notes, rebuilt its in-car AI compiler and runtime on MLIR and claims a 20 percent reaction-time improvement from the new stack, which, whatever discount you apply to vendor claims, is a statement about MLIR’s fitness for safety-critical, latency-bounded deployment that no benchmark suite could make.
Every one of these organizations faced the same choice: a decade of compiler archaeology from scratch, or a running start on shared infrastructure. None of them chose archaeology.
And beyond AI entirely, the quieter colonization: Flang, LLVM’s Fortran frontend, is built on MLIR dialects; CIRCT applies the same infrastructure to hardware design and verification; ClangIR is the upstream effort to give C and C++ themselves a structured MLIR layer above LLVM IR.
The pattern across all of it is the one Lattner himself named in January’s “Democratizing AI Compute, Part 8”: MLIR the infrastructure won more or less totally, while the original dream of one shared AI compiler on top of it fractured into vigorous, incompatible ecosystems.
The waist is shared; the programs flowing through it are not. Which is precisely what makes the waist worth money.
The economics of the waist
Regular readers know this publication’s core lens: in accelerated computing, margins pool wherever software friction is highest.
CUDA’s twenty-year lesson is that the compiler and kernel layer is not a cost center; it is the tollbooth. MLIR changes the tollbooth’s economics in three distinct ways, and the June acquisition finally put a market price on one of them.
First, it collapsed the entry ticket. Before this infrastructure existed, a credible accelerator software stack meant hundreds of compiler engineers for the better part of a decade; that is what CUDA cost, and it is why so many silicon startups died with excellent chips and unusable toolchains.
Previous sections showed the mechanism by which that changed: parsing, printing, verification, pass management, testing culture, all amortized across the industry, leaving a new entrant to spend only on what is genuinely theirs, the semantics of their silicon.
The evidence is the roster in Figure 4. Tenstorrent, a company of startup scale, fields a full graph-to-silicon compiler. AMD stands up an NPU stack as a side effort. This does not guarantee anyone beats NVIDIA; it guarantees the attempt no longer costs a billion dollars, and lowering the cost of attempts is how oligopolies erode.
Second, it standardized the engineers, not just the code. A compiler developer who learns dialects, patterns, and the linalg descent is productive at Google, AMD, Qualcomm, Tenstorrent, Modular, or any of a hundred teams, because they all speak the same infrastructure.
Before MLIR, compiler talent was siloed by proprietary IR; now there is a labor market for the waist. Every technology that has standardized its practitioners, from TCP/IP to Kubernetes, has seen the layer above it commoditize and the layer below it consolidate. Keep that in mind when reading the next paragraph.
Third, it created a control point you can buy. Qualcomm did not pay approximately $3.9 billion, per the Reuters-reported terms, for Modular’s revenue. It paid for roughly 150 people, per NAND Research’s headcount estimate, who include the founding architects of both MLIR and LLVM, plus MAX and Mojo, the most complete independent attempt to unify the layer above everyone’s silicon.
That is on the order of $26 million per employee, ten times acquihire norms, for a company whose language hit 1.0 beta in May, per its own PyPI listing, and whose compiler remains closed source with an open sourcing commitment for later this year.
The premium is not for what Modular sells; it is for what Qualcomm was structurally unable to build alone, demonstrated by the fact that its own Hexagon-MLIR effort was already consuming the ecosystem Modular’s founders created.
In the framework of our kernel-moat analysis: when the moat migrates from a proprietary language to a shared substrate, the scarce asset becomes the people who define the substrate’s direction, and June 24 was the day that asset got marked to market.
NVIDIA’s response is the most sophisticated play on the board: embrace the waist from above.
Contributing to Triton’s NVIDIA backends and shipping CuTe DSL on MLIR costs NVIDIA the exclusivity of its language layer, which was already eroding, while making its hardware the best-supported target of the shared stack everyone else depends on, and keeping the truly proprietary layers, ptxas, SASS scheduling, NVLink-domain libraries, exactly as closed as before.
Commoditize your complement, in textbook form. For the neoclouds and inference operators this publication serves, this is not abstract; it changes the arithmetic that turns capex into margin.
A fleet’s single largest source of pricing power against its accelerator vendor is credible multi-homing: the ability to qualify a second silicon supplier and actually move workloads.
Historically that ability died at the kernel layer, because a serving stack tuned to CUDA was a serving stack married to NVIDIA, and the cost of re-qualifying every fused kernel on a new target was prohibitive enough that “we could switch” was rarely true.
The shared waist is what makes the threat credible, incrementally. Every serving stack that targets the waist, vLLM and SGLang through Triton, IREE, MAX, is quietly reducing the switching cost between a B300 fleet and an MI400 fleet at the kernel layer.
Portability never arrives as a binary event, and CUDA-only paths (FlashAttention’s fastest branches, NCCL, the long tail of fused ops) remain very real. But price convergence per token happens at the margin, and the margin is now programmable.
Watch whether accelerator price-performance spreads narrow over the next four quarters as these stacks mature; that spread is the purest financial expression of what this article has been describing.
The counter-narrative
The strongest version of the case against, because a thesis untested against its critics is marketing.
“The library ate the compiler.” The performance frontier on NVIDIA hardware is still held by artisanal kernels: FlashAttention lineages, cuDNN, CUTLASS templates, hand-scheduled decode paths.
Much of what shipping compilers do is orchestrate calls into those libraries, and a fair reading of the TPP-MLIR work (arXiv 2404.15204), which reached roughly 90 percent of expert-written kernel performance on CPU benchmarks using upstream MLIR, is that the last ten percent is precisely where the money lives.
The rebuttal is not that compilers beat ninjas; it is that MLIR is the substrate the ninjas themselves are migrating to, CuTe DSL and Gluon being exhibits A and B. But the point stands: the waist wins by hosting excellence, not by automating it away.
“LLMs will write the kernels.” If models generate CUDA or PTX directly on demand, does a shared IR layer matter less? Plausibly it matters more: generated code needs verification, search needs a space with structure, and an IR with machine-checkable semantics and a transform language is a far better substrate for automated kernel search than free-form C++.
But intellectual honesty requires flagging this as the genuine wildcard: a world of cheap, correct, model-generated kernels reshuffles every layer of this analysis, ours included.
“Dialect soup.” Forty-eight upstream dialects, hundreds downstream, and no blessed end-to-end path is real fragmentation with real costs: two teams can both “use MLIR” and share nothing but a parser.
Lattner’s own Part 8 critique makes essentially this charge, and the upstream community’s response, charters, governance reform, the StableHLO-style versioning discipline spreading to more dialects, is work in progress, not a solved problem.
Relatedly, core MLIR offers no stable cross-release IR contract: downstream projects pay a real rebase tax every LLVM release, a tax Triton and IREE budget actual headcount for.
“The C++ tax.” The infrastructure’s expressiveness rides on heavy C++ templates and TableGen metaprogramming; build times are punishing, the learning curve is a cliff, and debugging a misapplied pattern through the greedy driver is an acquired skill.
Python bindings and tooling soften the edges for users, but dialect authors live in C++, and that gates the contributor pool. Fair, true, and priced in: the roster in Figure 4 is the market’s judgment that the tax is worth the leverage.
Every one of these criticisms describes a middle-aged infrastructure project with genuine adoption. None of them describes an alternative, and in infrastructure, the alternative is the only criticism that kills.
What to do with this
If this article did its job, MLIR stopped being a logo on other people’s architecture slides and became legible machinery.
Here is the shortest path from legible to useful, calibrated for readers who already know what a memory coalescing problem feels like.
Learn to read dumps before writing anything. The single highest-leverage habit is inspecting the IR your existing tools already produce. Triton will show you its TTIR and TTGIR for any kernel you own; start with a matmul you understand and find the layout encodings.
On your own experiments, mlir-opt --mlir-print-ir-after-all is the x-ray: every pass, before and after. When output confuses you, --mlir-print-op-generic strips the sugar and shows you the honest structure, exactly as in section 3.
Replay this article. Everything here ran on stock Ubuntu 24.04 packages: apt install mlir-20-tools llvm-20 libmlir-20-dev and you have mlir-opt, mlir-runner, mlir-translate, mlir-tblgen, and mlir-reduce (the last being delta-debugging for IR: feed it a crashing module and a script, get back a minimal reproducer).
The dossier below lists every command. An afternoon of replaying the walkthrough will teach you more than a month of architecture diagrams.
Then climb the same ladder the ecosystem climbed. The canonical study sequence: the Toy tutorial on mlir.llvm.org for the object model; the linalg dialect rationale for the structured-ops philosophy of section 5; the transform dialect tutorial for section 7’s machinery; then the CGO 2021 paper, which reads completely differently once you have touched the IR.
For a working CUDA engineer, a realistic 30-day arc is: week one, read IR and replay the walkthrough; week two, lower your own toy op through the same pipeline and break things on purpose; week three, build the previous section dialect and make the pattern fire on real IR; week four, open the dumps of one production Triton kernel you own and annotate every layout decision.
At the end of that month you will be conversant in the layer where, as section 10 argued, an increasing share of this industry’s margin gets decided.
Seven years ago this was a whiteboard sketch about taming TensorFlow’s compiler zoo. Today it compiles the kernels NVIDIA brags with, the NPU stacks of three chip giants, the vision system in a million cars, and it just priced a 150 person team at nearly four billion dollars.
The previous installment of this series ended by observing that LLVM won by being the infrastructure nobody had an incentive to replace. MLIR is running the same play one level up, on the layer where the AI hardware war is actually fought.
The empire had one IR. The dominion has as many as you can define, verify, and lower, and now you know how it is done.
Verification dossier
House methodology: every load-bearing claim is tiered, every number is sourced or explicitly derived, and for this piece, every code artifact was executed against real tools before publication.
Environment: Ubuntu 24.04, stock packages mlir-20-tools, llvm-20, libmlir-20-dev, all LLVM/MLIR 20.1.2.
Syntax note: MLIR evolves between LLVM releases; snippets are guaranteed against 20.1.2 and may need small spelling changes on older or newer toolchains.
Code verification manifest
Claims and confidence
Tiers: A primary source or machine-verified. B reputable secondary reporting or vendor primary for vendor facts. C reasoned inference or derived arithmetic, flagged as such in text. D interpretation and forecast: our analysis, argued not asserted.
Sources
Qualcomm, “Qualcomm to Acquire Modular,” investor press release, June 24, 2026. investor.qualcomm.com/news-events/press-releases/news-details/2026/Qualcomm-to-Acquire-Modular/default.aspx
Modular, “Qualcomm to acquire Modular,” company blog, June 2026. modular.com/blog/qualcomm-to-acquire-modular
Quartz, deal report citing Reuters and WSJ terms, June 24, 2026. qz.com/qualcomm-acquires-modular-ai-software-stock-deal-062426
NAND Research, “Qualcomm acquires Modular for its hardware-agnostic AI software layer,” June 2026. nand-research.com
Modular repository newsline (funding history, releases). github.com/modular/modular
Chris Lattner, “Democratizing AI Compute, Part 8: What about the MLIR compiler infrastructure?”, Modular blog, January 2026. modular.com/blog/democratizing-ai-compute-part-8-what-about-the-mlir-compiler-infrastructure
Modular, “The Path to Mojo 1.0”; PyPI, mojo-compiler release history (1.0.0b1, May 7, 2026). modular.com/blog/the-path-to-mojo-1-0; pypi.org/project/mojo-compiler
NVIDIA, CUTLASS documentation: CuTe DSL introduction and overview. docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_general/dsl_introduction.html
Triton project, Gluon documentation and dialect sources. triton-lang.org/main/gluon
Qualcomm et al., “Hexagon-MLIR,” arXiv 2602.19762, 2026. arxiv.org/pdf/2602.19762
Phoronix, “AMD MLIR-AIE 1.2,” January 2026. phoronix.com/news/AMD-MLIR-AIE-1.2
Tesla Oracle, “Tesla rolls out FSD v14.3 (2026.2.9.6): better reaction time, rewritten AI compiler with MLIR,” April 8, 2026. teslaoracle.com
Lattner, Amini, Bondhugula, Cohen, Davis, Pienaar, Riddle, Shpeisman, Vasilache, Zinenko, “MLIR: Scaling Compiler Infrastructure for Domain Specific Computation,” CGO 2021.
“TPP-MLIR” upstream performance study, arXiv 2404.15204. arxiv.org/pdf/2404.15204
MLIR project documentation: Toy tutorial, linalg rationale, transform dialect tutorial. mlir.llvm.org
StableHLO specification and repository. github.com/openxla/stablehlo
Tenstorrent, tt-mlir repository. github.com/tenstorrent/tt-mlir
Every time your compiler converts a tensor operation into flat memory pointers, it destroys the exact geometric facts needed to keep hardware running fast. Traditional compiler pipelines spend immense compute trying to reconstruct loop nests and aliasing bounds from opaque pointers. MLIR flips this entirely by keeping structure intact across nested dialect layers. 🛠️
When you lower a fused matrix multiply with bias and ReLU down to hardware PTX, the crossing from value tensors to mutable memrefs is where performance dies or survives. Tensor algebra guarantees zero aliasing. The moment you strip that away, the backend must assume any reader can observe intermediate sums, forcing redundant global HBM stores on every reduction step. That's why destination-passing style matters so much. By declaring output allocations back in pure value land, one-shot bufferization plans in-place memory updates before a single physical address gets assigned. 💡
What if we pushed those non-aliasing tensor proofs directly into hardware crossbar gates? Instead of letting lowering pipelines discard structural guarantees, routing value-level SSA proofs straight into local SRAM clock-enable registers would freeze un-gated bus polling during matrix solves. That eliminates HBM writeback penalties entirely at the silicon boundary. ⚡
Are your custom kernel passes holding onto tensor-level non-aliasing facts through bufferization, or is your lowering pipeline silently paying a memory tax on every multiply? 🤨