What Mojo's Optimiser Leaves Out

A first look at the compiler and some learnings about Mojo's design choices

Ever since Chris Lattner founded Modular and announced Mojo as a new programming language with the ease-of-use of Python but the performance of systems programming languages like C++, I have been interested. So much so that I applied to one of the first openings at Modular. Unfortunately, openings at Modular were limited to the US.

For the last three years I have researched relaxed memory models: how you reason about concurrent programs when the compiler and the processor are both allowed to reorder your instructions, and how you ensure that the optimisations a compiler performs are sound with respect to the memory consistency model of the programming language and the hardware architecture. Now that Modular has open-sourced Mojo^[1], the first thing I wanted to see was which optimisations Mojo performs in order to make good on that famous claim of systems-language-level performance. The observations below are the result of a cursory read through the source, and of compiling a batch of small test programs to inspect the IR and the assembly that come out.

Some background first. Mojo reached 1.0 on 11 August, and a week later Modular – since July part of Qualcomm – published the compiler, the toolchain and everything needed to build them under an Apache-2.0 license with LLVM exceptions. For the four years before that the standard library was open and the compiler was a binary. So this week is the first time anyone outside the company has been able to look at how Mojo actually generates code. Since it compiles to LLVM, I expected the typical LLVM frontend, but that is not exactly what is going on.

The pipeline is a fork, not a call

The compiler is historically code-named KGEN (kernel generator). It is an MLIR stack^[2]: six dialects, its own optimiser at the MLIR level, and LLVM underneath. It is one part of a repository that also holds the Mojo language and its standard library, the AsyncRT runtime, and MAX – the kernel library and graph runtime built on top of all of it. What follows is about the compiler.

Most LLVM frontends call the LLVM library to build a pipeline with PassBuilder::buildPerModuleDefaultPipeline, and the library will construct a pipeline configured with defaults for the given optimisation level, -O3. KGEN constructs the pipeline by hand in LLVMPassesPipeline.cpp, and that file is recognisably a copy of upstream’s PassBuilderPipelines.cpp – the same function names (buildFunctionSimplificationPipeline, addInlinerPasses, addVectorPasses), and the same optimisation pass order. Even upstream’s comments and FIXMEs are carried across verbatim.

Interestingly, two passes are missing:

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13

static void addVectorPasses(FunctionPassManager &FPM,
                            const CompilationOptions &options) {
  FPM.addPass(LoopLoadEliminationPass());
  FPM.addPass(InstCombinePass());
  FPM.addPass(SimplifyCFGPass(...));
  FPM.addPass(SLPVectorizerPass());       // SLP, but no LoopVectorize
  FPM.addPass(VectorCombinePass());
  FPM.addPass(InstCombinePass());
  // "Now that we are done with loop unrolling, be it either by LoopVectorizer,
  //  or LoopUnroll passes, ..."          <- upstream comment, now stale
  FPM.addPass(SROAPass(SROAOptions::PreserveCFG));
  ...
}

LoopVectorizePass and LoopUnrollPass appear nowhere in the compiler – not under a flag, not behind a target hook, nowhere in the tree. The upstream comment that still refers to them is the fingerprint of a deletion rather than a divergence.

The vectoriser’s absence, at least, is not news. Two months before any of this source was published, someone asked on Modular’s forum why their generated DSP code would not vectorise the way the equivalent C++ did; an engineer replied that Mojo “intentionally disable’s LLVM’s autovectorizer"^[3], and gave the reasoning below.

Moreover, the unroller went with the vectoriser. The removal was done by forking the pipeline rather than configuring it, which adds long-term maintenance overhead and shows intent. And – further down – a great deal was left in. None of that is in the compiler documentation, which covers this stage in one line^[4]: “LLVM Optimization: Run LLVM’s optimization pipeline”.

The straight-line vectoriser stays. SLPVectorizerPass and VectorCombinePass are both scheduled, so adjacent scalar operations are still packed and existing vector code is still cleaned up. What is gone is specifically loop vectorisation and loop unrolling.

What that looks like

A textbook AXPY:

1
2
3
4

def axpy(a: Float32, x: Pointer[Float32, MutAnyOrigin],
         y: Pointer[Float32, MutAnyOrigin], n: Int):
    for i in range(n):
        y.unsafe_store(i, a * x.unsafe_load(i) + y.unsafe_load(i))

On a 13th-generation Core i7 with AVX2 and fused multiply-add (FMA), at -O3, that is the entire function:

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15

axpy_f32:
	xorl	%eax, %eax
	testq	%rdx, %rdx
	cmovgq	%rdx, %rax
	jle	.LBB0_3
	xorl	%ecx, %ecx
.LBB0_2:
	vmovss	(%rdi,%rcx,4), %xmm1
	vfmadd213ss	(%rsi,%rcx,4), %xmm0, %xmm1
	vmovss	%xmm1, (%rsi,%rcx,4)
	incq	%rcx
	cmpq	%rcx, %rax
	jne	.LBB0_2
.LBB0_3:
	retq

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20

; main vector loop only; the full function also has
; a runtime alias check, an 8x-unrolled scalar path
; and a scalar remainder loop
	vbroadcastss	%xmm0, %ymm1
.LBB0_10:
	vmovups	(%rdi,%rcx,4), %ymm2
	vmovups	32(%rdi,%rcx,4), %ymm3
	vmovups	64(%rdi,%rcx,4), %ymm4
	vmovups	96(%rdi,%rcx,4), %ymm5
	vfmadd213ps	(%rsi,%rcx,4), %ymm1, %ymm2
	vfmadd213ps	32(%rsi,%rcx,4), %ymm1, %ymm3
	vfmadd213ps	64(%rsi,%rcx,4), %ymm1, %ymm4
	vfmadd213ps	96(%rsi,%rcx,4), %ymm1, %ymm5
	vmovups	%ymm2, (%rsi,%rcx,4)
	vmovups	%ymm3, 32(%rsi,%rcx,4)
	vmovups	%ymm4, 64(%rsi,%rcx,4)
	vmovups	%ymm5, 96(%rsi,%rcx,4)
	addq	$32, %rcx
	cmpq	%rcx, %rax
	jne	.LBB0_10

Scalar, not unrolled, and no runtime alias check – because nothing downstream was going to need one. The same loop written in C and put through clang comes out as a 256-bit main loop unrolled four times, a runtime disambiguation check to establish that x and y do not overlap, an eight-times-unrolled scalar path, and a scalar remainder.

The vfmadd213ss in the Mojo output is worth noting separately: FMA contraction did happen. KGEN attaches the contract fast-math flag by default, so the IR it emits carries fmul contract and fadd contract without anyone asking for -ffast-math. Only contract, though – not reassoc, not nnan, not ninf.

Because SIMD is in the type system

I read the missing auto-vectorisation pass as a deliberate move: away from code that is implicitly vectorisable at compile time, and towards code that is explicitly vectorised through the language’s own portable SIMD library^[3]. Discovering vectorisation is not a job Mojo has taken on.

In KGEN’s IR, a SIMD vector is a first-class parameterised type:

1
2
3
4
5
6
7

def KGEN_SIMDType : KGENType<"SIMD", "simd", [...]> {
  let summary = "parameterized SIMD vector type";
  let parameters = (ins "TypedAttr":$size, "TypedAttr":$dType);
  ...
  std::optional<int64_t> getResolvedSize() const;
  bool isScalar() const { return getResolvedSize() == 1; }
}

The width and the element type are TypedAttr – parameter expressions, not integers. simd<W, f32> with a symbolic W is a well-formed type; that is why getResolvedSize() returns an optional. A scalar is just simd<1, T>, which is why in Mojo Float32 really is SIMD[float32, 1] and why the same unsafe_load[width=W] works for both.

The vector width is declared in code, preserved by the front-end, and written into the LLVM IR:

1
2
3
4
5
6

comptime W = simd_width_of[DType.float32]()   # 8 here

for i in range(0, n, W):
    var xv = x.unsafe_load[width=W](i)
    var yv = y.unsafe_load[width=W](i)
    y.unsafe_store[width=W](i, a * xv + yv)

1
2
3
4
5
6
7
8

	vbroadcastss	%xmm0, %ymm0
.LBB0_2:
	vmovups	(%rdi,%rax,4), %ymm1
	vfmadd213ps	(%rsi,%rax,4), %ymm0, %ymm1
	vmovups	%ymm1, (%rsi,%rax,4)
	addq	$8, %rax
	cmpq	%rdx, %rax
	jl	.LBB0_2

The scalar a has to be copied into all eight lanes before it can multiply a vector. That broadcast is a splat, and it is the vbroadcastss sitting above the loop label. It is already in the IR before any LLVM pass runs – an insertelement putting a into lane zero, then a shufflevector copying lane zero across the rest of the <8 x float> – both emitted by the frontend. LLVM’s remaining job is to select the instruction and hoist it out of the loop. Vectorisation was never a discovery problem.

What it costs

I measured a polynomial reduction over four million float32, twenty repetitions:

5.6× of an available 8×, for writing the width yourself. On a compiler that auto-vectorised, the first two rows would be the same number.

The trade is clearly predictability. Modular have said so: asked on their forum in June why a DSP kernel was not being vectorised, Owen Hilyard replied^[3] that “Mojo intentionally disable’s LLVM’s autovectorizer, since Mojo ships with a portable SIMD library in the stdlib”, and that this “makes SIMD code far less brittle and often more portable” – adding that in compute kernels “the autovectorizer can actually get in the way of high performance code by vectorizing your scalar path when your vector ALUs are already occupied”.

Auto-vectorisation is the largest single source of unpredictable performance in C and C++. It fires or it does not, depending on alias analysis, trip-count analysis, and cost models that shift between compiler releases; a kernel that vectorised last year may quietly stop. For a language whose entire pitch is that you can see the machine from the source, refusing to guess is a coherent choice. It does mean that naive Mojo will lose to C on scalar loops, and only wins once the SIMD is written out.

Is any of that unusual?

Putting the vector width in the type is close to consensus. It has been arrived at repeatedly and independently:

So the type itself is the ordinary part. Two details are less so:

The first is that the width is a parameter expression rather than an integer. In most of the table the length is a constant by the time the type exists – a const generic, a template argument, a literal. In KGEN it is a TypedAttr, so simd<W, f32> with a symbolic W is a well-formed type in the compiler’s own IR, carried through optimisation and resolved during elaboration. That is why getResolvedSize() returns an optional rather than an integer. The nearest relatives here are the scalable vector extensions: an Arm SVE svfloat32_t is a type whose length genuinely is not known when you write it. Mojo’s width will be known, but not yet – which is a different thing again, and closer to how a dependently-typed array length behaves than to a template parameter.

The second is that scalar and vector are the same thing. A scalar is simd<1, T>; isScalar() is literally getResolvedSize() == 1; Float32 in Mojo really is SIMD[float32, 1]. That is why one unsafe_load[width=W] serves both, and why generic code can be written once and instantiated at either.

What distinguishes Mojo here is not really the type system. Look at the last column of tbl2. Zig, Rust and C++ all give you portable vector types and a loop vectoriser underneath: explicit SIMD when you want control, automatic vectorisation when you do not. They are not alternatives, they are layers. Mojo is the entry in that table that removed the lower layer.

The nearest relative is ISPC, whose premise is that you write the parallelism and the compiler does not guess – and which Modular were evidently thinking about, since KGEN’s own design overview says “I’m particularly interested in exposing an ISPC-style SPMD transformation as a generator”. The lineage is not hidden.

For comparison, portable SIMD is not available on stable Rust at all: it remains nightly-only, so the realistic stable-Rust options are architecture intrinsics or trusting the vectoriser. C++26 has only just standardised its version. Mojo has shipped the explicit route as the only route, in 1.0, in a language whose users are largely writing numerical kernels.

What Mojo kept

The memory consistency model of a programming language, such as Mojo, has to account for the optimisations that move memory accesses relative to one another, and recent work models them explicitly^[10]: store-to-load forwarding, dead-store elimination, constant propagation, and hoisting an access out of a conditional branch. Every one of them is still scheduled – see tbl3:

As a researcher in concurrent algorithms I went looking for concurrency and atomics first, which is why that list caught my eye. But I would not read the passes that survived as a judgement about concurrency, because I do not think anyone was judging. Modular’s stated vision is heterogeneous compute – “the one programming language developers can use to target all the diverse hardware"^[11] – not shared-memory synchronisation. Mojo ships no memory model document. First-class async “fully integrated with Mojo’s type and memory models” is still an open item on the roadmap. parallelize is not in the standard library at all; it lives in MAX. Concurrency is plumbing here, not a subject.

LoopDistributePass is still scheduled – a pass whose only purpose is isolating dependencies so that partial loops can vectorise, and which has nothing to feed. So is InjectTLIMappings, which populates scalar-to-vector function mappings. The comments quoted earlier still describe cleanup after a pass that no longer runs. Two addPass lines came out of a copied file and nothing around the cut was tidied.

It is worth adding what the memory model itself looks like, since it is the question I went in with. Mojo’s atomics are LLVM’s atomics, unchanged: the ordering lattice in its POP (Parametric Operations) dialect is llvm::AtomicOrdering with the same names and the same numbering, and the standard library re-exports it more or less directly. Every atomic operation additionally carries a syncscope, which Mojo makes a compile-time parameter of the atomic’s type rather than an argument at each use. In the zoo it fits with the LLVM memory model, with a scope dimension added. Defaults are sequentially consistent.

Odds and ends

A few other things I noticed.

• Unrolling is explicit too. The LLVM unroller is gone, but Mojo has had an @unroll decorator and MLIR-level LoopUnrolling pass for years^[12] – the same substitution as with SIMD: a general-purpose pass replaced by something the programmer asks for by name.
• There are only two optimisation levels. getOptimizationLevel maps 0 to O0 and everything else to O3. -O1 and -O2 are accepted on the command line and both silently mean -O3. There is no Os or Oz path.
• LLVM is built from source and patched, but the optimiser is stock. Six patches are applied. All six are build plumbing or LLDB bug fixes – a symbol export list, a Bazel select() revert, a heap-corruption fix in ObjectFileELF, an ORC header cycle. None touches the optimiser, the vectoriser, alias analysis or the memory model. Whatever Mojo does differently, it does not do by patching LLVM.
• KGEN reimplements the classical mid-end at MLIR level. Mem2Reg, SROA, SCCP, SimplifyCFG, an inliner, argument promotion, loop unrolling – all of them again, on MLIR. This is because the compile-time parameter system has to be constant-folded, specialised and inlined away before the program can be lowered to LLVM at all, and LLVM has no representation for any of it. Once you need SCCP and an inliner up there anyway, the rest is close to free.
• The async runtime uses performance cores only. The one runtime fact in this list rather than a compiler one – AsyncRT sits beside the compiler in the repository, and the standard library binds to it. Workers are pinned by affinity, not merely counted. E-cores and SMT siblings go unused by default.

How Mojo achieves its performance

So where does the systems-language performance actually come from? Four things, none of them exotic.

LLVM, doing what it has always done. There is no new optimisation technology here. KGEN’s job is to resolve the parameter system and hand LLVM ordinary monomorphic IR, and it hands over good IR: getelementptr inbounds, noundef on parameters, fastmath contract already attached. From there it is upstream’s -O3, near enough. A compiled, statically typed language handing well-formed IR to LLVM is most of the answer to “why is it as fast as C++”, and it is the least surprising part of the story.

Predictability first. The two passes that were taken out are the two whose behaviour is hardest to predict from the source: whether a loop vectorises and by how much it unrolls depend on alias analysis, trip counts, and cost models that move between releases. Removing them costs peak throughput on naive code and buys the guarantee that the machine code has the shape of the loop you wrote.

Explicit parallelism in their place. What was removed is not absent, it is relocated into the language, where you have to be explicit about it. Vector width is a type parameter, vectorize handles the drain loop, @unroll is a decorator, parallelize is a library call. The transformations other compilers try to infer are things Mojo makes you say.

And everything else left alone. SLP vectorisation, GVN, LICM, DSE, SCCP, inlining, FMA contraction by default – the ordinary machinery is all still scheduled, much of it twice, once on MLIR and once on LLVM IR. Scalar Mojo is compiled competently. It simply is not vectorised for you.

So the claim holds, with a condition attached. Python’s ergonomics, yes. A systems language’s speed – when you write like a systems programmer. The 5.6× in tbl1 is the size of that condition, and on a loop this ordinary a C compiler would have closed it for you. What Mojo offers is not a compiler that goes looking for the performance on your behalf; it is a language where the thing you would otherwise have dropped to intrinsics for is a type parameter instead.

Two things worth watching

A forked pipeline is a standing cost. Copying hundreds of lines of upstream means owning them: every time LLVM’s -O3 changes, somebody has to decide whether to follow. Configuration hooks would have removed two passes at a fraction of the price. Carrying the fork only makes sense if you mean to keep diverging – and the scaffolding for that is already in place. There are three Modular-written LLVM IR passes in the tree, a TargetBackend interface with a hook for splicing target-specific passes into the pipeline, a run of commits moving backend-specific logic out of the shared pipeline and into per-target backends, and an opt-alike driver of their own to run any of it in isolation. I would read the fork as a bet on more divergence from stock LLVM, not less. Two deletions is where it starts, not where it ends.

Concurrency is still in its infancy. For a language at 1.0 the gap is striking. There is no memory model document. First-class async, “fully integrated with Mojo’s type and memory models”, is an open item on the roadmap. parallelize is not in the standard library. The atomics are LLVM’s, re-exported with the names and numbering unchanged, which means Mojo’s memory semantics are for now whatever LLVM IR happens to mean – a reasonable place to start, but not the last word. Given that the stated vision is heterogeneous compute, none of this is surprising. It is the half of the language I will be watching.

Reproducing this

Everything above comes from the public repository^[13] and the shipped compiler. mojo build will show you both ends of the pipeline:

1
2

mojo build --emit llvm -O 3 f.mojo -o f.ll   # IR after KGEN, before LLVM
mojo build --emit asm  -O 3 f.mojo -o f.s    # after the full LLVM pipeline

The CLI calls --emit llvm “unoptimized LLVM IR”, which is true only of LLVM’s passes – KGEN’s own MLIR pipeline has already run, and -O0 versus -O3 change that output substantially. Diffing the two isolates the frontend’s optimiser; diffing --emit llvm against --emit asm isolates LLVM’s contribution.

References

1. Modular Inc.: Mojo is now open source, 18 August 2026 www.modular.com/blog/mojo-open-source
2. Lattner, Amini, Bondhugula, Cohen, Davis, Pienaar, Riddle, Shpeisman, Vasilache, Zinenko: MLIR: Scaling Compiler Infrastructure for Domain Specific Computation, CGO 2021 10.1109/CGO51591.2021.9370308
3. Owen Hilyard (Modular): reply in Mojo auto-vectorization for generated DSP code, Modular forum, 19 June 2026 forum.modular.com/t/mojo-auto-vectorization-for-generated-dsp-code/3252
4. Modular Inc.: Mojo compiler walkthrough, KGEN/docs/MojoCompilerWalkthrough.md, “Phase 6: Lowering to LLVM”. Permalink pinned at commit 33cd4694 github.com/modular/modular/blob/33cd4694b19649bec7f5acac88b0430371805dc6/KGEN/docs/MojoCompilerWalkthrough.md#L561
5. Zig Language Reference: Vectors (@Vector(N, T)) ziglang.org/documentation/master/#Vectors
6. rust-lang/portable-simd: The testing ground for the future of portable SIMD in Rust. Nightly-only as of August 2026 github.com/rust-lang/portable-simd
7. Data-parallel types (SIMD), C++26 (P1928, merged from Parallelism TS 2) en.cppreference.com/w/cpp/numeric/simd
8. Arm: Introduction to SVE — sizeless vector types (svfloat32_t) whose length is not a compile-time constant developer.arm.com/documentation/102131/latest/
9. Pharr, Mark: ispc: A SPMD Compiler for High-Performance CPU Programming, InPar 2012. uniform/varying type qualifiers ispc.github.io/
10. Richards, Wright, Cooksey, Batty: Symbolic MRD: Dynamic Memory, Undefined Behaviour, and Extrinsic Choice, OOPSLA 2025 2025.splashcon.org/details/OOPSLA/104/Symbolic-MRD-Dynamic-Memory-Undefined-Behaviour-and-Extrinsic-Choice
11. Modular Inc.: Mojo vision — “the one programming language developers can use to target all the diverse hardware—CPUs, GPUs, custom accelerators, ASICs, and more” docs.modular.com/mojo/vision/
12. Modular Inc.: What is loop unrolling? How you can speed up Mojo code with @unroll www.modular.com/blog/what-is-loop-unrolling-how-you-can-speed-up-mojo
13. Modular Inc.: modular/modular, GitHub. Read at commit 33cd4694 on main and at tag mojo/v1.0.0 github.com/modular/modular