Introduction
Four days after Modular shipped Mojo 1.1, I did what I just always do with a brand new compiler. I installed it in a container with one CPU core and no GPU, wrote a kernel that doubles a vector, and asked for PTX for an H100. I asked for sm_90. The first two directives of what came back were these:
.version 8.5
.target sm_90aThat one letter is the subject of this piece. sm_90 is Hopper. sm_90a is Hopper with its architecture-specific instructions unlocked, among them the warpgroup matrix multiply and the warpgroup register reallocation that the fastest Hopper kernels are built on.
The price of the suffix is written into NVIDIA’s own assembler. PTX that declares .target sm_90a can only be compiled for sm_90a. It cannot be compiled for plain sm_90, and it cannot be carried forward to the next generation. I didn’t ask for that trade.
The compiler made it for me, without a warning, and it makes the same trade for every Hopper and Blackwell name it accepts but one.
I spent the following days pulling on that thread, and it runs through the whole stack. I installed Mojo 1.1.0 and MAX’s max-core 26.6.0 from PyPI, cloned the compiler Modular open-sourced in August, took NVIDIA’s ptxas 13.4 from PyPI to check what Mojo emits, and compiled well over a hundred small probe programs.
Nothing here ran on a GPU. Everything is about what the compiler decides before a GPU is involved. I’m not a compiler engineer, I’m someone who likes taking tools apart, and Mojo is now a tool you can take apart all the way down. Where I infer instead of measure, I say so, and the dossier at the end grades every claim.
The short version is basically this: Mojo’s portability is real, and it lives entirely in the source. The compiler treats a target name as a product rather than an instruction set and compiles for the exact chip it is told about. When the compiler travels with the program, as it does inside MAX’s runtime library, that costs nothing.
When it does not, as in an executable made with mojo build, the kernels inside are locked to one chip: an executable built for an H100 carries PTX that NVIDIA’s own assembler refuses to build for any Blackwell part. Either way, every NVIDIA kernel is finished by NVIDIA’s closed assembler, which MAX ships inside its wheel.
The design is, at least in my own view, pretty coherent. In several places the implementation is not: a device table that quietly compiles a Jetson Orin as an A100, a Blackwell guard that reads the machine you compile on instead of the machine you compile for, and a Jetson Thor entry that cannot reach the tensor memory its hardware has.
What changed since the Mojo you remember
The Mojo of 2023 was (probably) pitched as a superset of Python with a split personality. def behaved like Python: dynamic, permissive, allowed to raise. fn was the strict twin: typed, non-raising unless declared, arguments immutable by default.
That split was the idea people remembered. It is gone. Modular said in August 2025 that Mojo may or may not grow into a full Python superset and that this was acceptable, which is how Simon Willison summarized the change when the compiler went open.
The repository’s own release notes date the rest. fn appears in the November 2022 notes as a new, stricter declaration. It is deprecated in 0.26.2, warned on by the compiler from 1.0.0b1 in May 2026, and finally removed in 1.1.0, released on 17 September 2026, together with alias, the __comptime_assert keyword, and the @parameter if and @parameter for forms.
The 1.0 announcement promised that changes during 1.x would be primarily additive, with breaking changes managed carefully in the way mature languages like C++ manage them. The first minor release after 1.0 deleted three keywords and two syntax forms.
To be fair, all of them were deprecated during the 1.0 cycle, so code that compiled without deprecation warnings under 1.0 should still compile under 1.1. The more interesting thing is what the removals leave behind. We wrote a batch of small programs to test what 1.1 actually enforces, and ran each one through the compiler:
def now means what fn used to mean. It does not raise unless it says raises, every variable is introduced with var, and arguments default to the immutable imm convention.
The memory model underneath is affine by default: a value moved out with ^ is uninitialized, and using it is a compile error. It is linear on request: a type declared with Deinitable where False never conforms to Deinitable, so it must be explicitly consumed, and the compiler rejects any path that drops it.
Rust has affine types and no native linear ones; Mojo has both in one type system. Exclusivity is enforced at call sites, and since 1.0 the lifetime checker, experimentally, understands the inside of containers, so a reference into a List held across an append is rejected instead of dangling after reallocation. Strings refuse to have a length until you say which length you mean.
What is still missing is also clear. Pattern matching exists on the main branch as an experimental __match statement, and the nightly changelog notes that exhaustiveness is not yet checked for enums or Bool. An async model and unions are on the roadmap.
So Mojo 1.1 is not Python with types anymore. It reads like a Rust whose surface borrowed Python’s indentation, attached to a compile-time metaprogramming system more capable than either. We have argued before that strict semantics and one way to do things suit a world where models write a growing share of the code.
Mojo 1.1 reads like a language that took that argument seriously, whether or not anyone at Modular would phrase it that way.
What pip actually installs
pip install mojo pulls five packages: mojo itself (a 48.9 MB wheel holding the language server, the debugger and the REPL entry point), mojo-compiler (an 80.7 MB wheel) and its companion mojo-compiler-mojo-libs, mblack 26.6.0 for formatting, and mojo-lldb-libs.
The directory it installs is 675 MB. The compiler driver is a single binary of 142,215,032 bytes. It carries its own linker, an lld of 115,567,904 bytes, which is 81% of the size of the compiler that calls it. The standard library ships as one precompiled container, std.mojoc, 3,226,898 bytes, with the magic bytes MPKG and an entropy of 7.999 bits per byte.
Its entropy says it is compressed, or encrypted, and either way it is opaque. The source of that library is on GitHub under Apache 2.0; the file the compiler actually loads is not something you can read.
A “hello world” builds in 2.33 seconds cold and 0.60 seconds from a warm cache in our first session, and in 2.97 to 3.03 and 0.74 to 0.76 seconds across three later runs on the same, busier shared core. The ratio held: a warm build costs a quarter of a cold one.
The result is a 17,944-byte executable. The executable is not standalone. It links libKGENCompilerRTShared.so through an absolute RUNPATH that points into the Python environment Mojo was installed into, /usr/local/lib/python3.12/dist-packages/modular/lib in our case. Move the binary to a machine without that directory and the loader will not find its runtime unless you point it there. That is a normal trade for a young toolchain, and worth knowing before you ship anything built this way.
The bigger surprise is what the wheel does not contain. Mojo is sold as the language for GPUs, and after pip install mojo there is no GPU package. from gpu.id import block_idx fails with unable to locate module 'gpu'. The 1.0 changelog explains it: the accelerator APIs moved into a new max Mojo package and the layout package was bundled with MAX instead of Mojo.
Install max-core 26.6.0, an 85.2 MB wheel, and a max package appears with max.gpu inside it, one of 31 precompiled packages totalling 24.4 MB: nn at 6.8 MB, linalg at 4.8 MB, layout, kv_cache, quantization, comm, shmem, state_space, and five vendor binding packages we come back to later.
It also brings libmax.so at 162.6 MB, the GPU runtime libMGPRT.so at 44.7 MB, and a 37.7 MB file called libNVPTX.so that turns out to be NVIDIA’s assembler. With both installed, the tree is 982 MB.
The licenses deserve one careful paragraph, because the headline from August was that Mojo is now open source, and that is true of the source. The PyPI metadata for mojo 1.1.0 and for max 26.6.0 both declare LicenseRef-MAX-Platform-Software-License.
The Python shim that provides the wheel’s entry points opens with a header calling itself Modular proprietary. The repository’s LICENSE is Apache 2.0 with LLVM exceptions.
Reporting from ModCon by NAND Research and RuntimeWire says Modular also rewrote the MAX license there to remove device restrictions and is moving MAX toward a source-available model. We are not lawyers and read none of this as a problem.
What is measurable is narrower: the tree you clone and the wheel you install carry different license identifiers, and if you need the Apache terms end to end, Modular’s own instructions build the compiler from source and run a file with ./bazelw run --config=build-mojo KGEN:mojo -- run hello.mojo.
The same post notes that a prebuilt compiler is still required today if you customize MAX kernels or models.
The compiler, now that it can be read
We cloned modular/modular at commit 26cfe94, dated 21 September 2026, and fetched the mojo/v1.1.0 tag (commit c6fa49f, 17 September) to compare against what actually shipped. The working tree is 11,054 files.
The compiler proper, the C++ and TableGen under Mojo/lib, Mojo/include and Mojo/tools, is 314,179 lines. The standard library, written in Mojo, is 185,284 lines. The compiler’s tests are another 149,354.
Then there is MAX: 1,048,124 lines of Mojo across kernels, models and tests, and 880,903 lines of Python. The Mojo code inside MAX alone is 3.3 times the size of the compiler. The language is the smaller asset in its own repository.
Under Mojo/include/Mojo there are five MLIR dialects, and their own descriptions say what each one is for. lit holds the language-level abstractions and lowers to kgen. pop defines parametric operations on top of MLIR’s LLVM dialect. kgen holds the top-level and support operations of what its description still calls the kernel generation framework. hlcf is a structured, high-level control-flow dialect, and co models async functions as coroutines with arbitrary suspension points.
Counting the TableGen definitions gives 196 operations, 114 types and 174 attributes across the five, and 45 passes declared next to them, plus 2 more in the shared Support library. The shape of kgen is the telling one: 39 operations, 64 types and 98 attributes, in 9,238 lines of TableGen, more than lit and pop combined.
Compile-time values in Mojo are carried as MLIR attributes; the device records we meet below are literally #kgen.target<...> attributes. That is why the dialect that carries a program’s compile-time structure is mostly types and attributes.
The name KGEN is everywhere once you look for it. Modular’s documented build target for the compiler is KGEN:mojo. The runtime every Mojo executable links is libKGENCompilerRTShared.so. The tools directory contains kgen, kgen-doc and kgen-reduce, and part of the test suite lives under test/kgen.
In our July piece on MLIR we described KGEN, from the public material available then, as a proprietary kernel generator that sat as a layer above Mojo. With the source open, that picture was wrong. KGEN is the compiler’s historical name and one of its dialects, sitting inside Mojo, not above it.
The directory list also has an Interpreter and an Elaborator. By their names and the passes they feed, the first runs comptime code inside the compiler and the second instantiates parametric code, which is where a large share of every build’s time goes.
Where the compile time goes
Mojo can time its own pipeline. --mlir-timing reports every MLIR pass, --llvm-timing every LLVM pass, and --timing-json writes both as one object. Mojo caches compiled modules in MODULAR_CACHE_DIR and a warm cache skips most of the pipeline, so every measurement here points it at an empty directory first.
For the “hello world” example, 46 distinct passes run, plus five analyses. The four largest are importing the Mojo modules at 16.6% of pipeline time, VerifyParameters at 16.1%, LowerLIT at 13.1% and InlineParametric at 12.8%, followed by CheckLifetimes, the borrow checker, at 6.8%.
For a program that compiles the vector kernel for sm_90a as an offload target, the order shifts: imports 19.3%, ElaborateGenerators 15.4%, VerifyParameters 12.7%, InlineParametric 9.4%, LowerLIT 8.6%, CheckLifetimes 4.7%.
The three passes that exist to instantiate and check parametric code take 31.7% of the hello-world pipeline and 37.5% of the GPU one. The tree view shows the elaborator working in rounds: InlineParametric followed by VerifyParameters, repeated, with dead-symbol elimination and canonicalization in between.
Then there is LLVM. In the timing run for the sm_90a program, which the LLVM timers force onto one thread, the MLIR pipeline took 5.92 seconds and every LLVM pass on both the x86 host and the NVPTX offload added up to 15.6 milliseconds, or 0.26% of the build. Instruction selection for the whole GPU kernel took less than half a millisecond.
In our piece on Triton, Triton’s own MLIR pipeline was 81% of compile time and ptxas the other 19%. In Mojo the assembler is not in the build at all, because PTX is turned into machine code when the program loads, and LLVM is a rounding error. Nearly everything the compiler spends its time on is deciding what your program means.
That made us expect compile time to grow with the number of specializations, since generating kernels instead of enumerating them is the whole bet behind the language.
So we built programs that instantiate a small parametric function work[N] once, 8, 32, 128 and 256 times through a comptime for, each from an empty cache. We ran each three times. The medians were 3.08, 3.09, 3.08, 3.18 and 3.21 seconds, and a least-squares line through them gives about 0.56 milliseconds per specialization on a floor of about 3.1 seconds: two hundred and fifty-six specializations add roughly 0.14 seconds.
At this size an extra specialization costs almost nothing. The floor is paid before any user code. An empty def main(): pass builds in 2.92 seconds cold, median of three, and its pipeline looks like hello world’s: importing modules 18.1%, VerifyParameters 16.7%. With no user code at all, that is the standard library being imported and its parametric code verified. Adding print(1) costs about 0.13 seconds more, print("hello") about 0.06.
The honest caveat is that our bodies were quite small. A kernel that unrolls deeply at compile time will cost more, and the compiler has a --loop-unrolling-warn-threshold flag, defaulting to 1024, for when it does.
The target sweep
mojo build --print-supported-accelerators lists every GPU the compiler will target. In 1.1.0 that is fourteen AMD architectures from gfx90a (MI250X) through CDNA3 and CDNA4 to RDNA2, RDNA3, RDNA3.5 and RDNA4 consumer parts, plus the mi300a APU name and two aliases.
Ten Apple entries, M1 through M5 each with and without Metal 4; and twenty NVIDIA names from sm_52 (Maxwell) to sm_121a (DGX Spark). --print-supported-targets lists the host CPU backends: the AArch64 family, 32 and 64-bit RISC-V in both endiannesses, and 32 and 64-bit x86. There is no Hexagon, no Qualcomm accelerator, no TPU and no Trainium in either list, which will matter at the end.
Modular’s GPU package exposes a function, _compile_code, that compiles a kernel for a named target at compile time and returns the assembly as a string. You don’t need to own any particular device.
We compiled the same small kernel for all twenty NVIDIA names and fed every result to NVIDIA’s ptxas 13.4.92, taken from the nvidia-cuda-nvcc wheel on PyPI.
from max.gpu.host.compile import _compile_code, get_gpu_target
from max.gpu import thread_idx, block_idx, block_dim
def vadd(c: Pointer[Float32, MutAnyOrigin], a: Pointer[Float32, ImmutAnyOrigin], n: Int):
var i = Int(block_idx.x * block_dim.x + thread_idx.x)
if i < n:
c[unsafe_offset=i] = a[unsafe_offset=i] * 2.0
def main() raises:
print(_compile_code[vadd, target = get_gpu_target["sm_90"]()]())Five names come back with a suffix nobody asked for: sm_90 becomes sm_90a, sm_100 becomes sm_100a, sm_103 becomes sm_103a, sm_120 becomes sm_120a and sm_121 becomes sm_121a.
One name goes the other way: sm_87, Jetson Orin, comes back as .target sm_80. And sm_110a, Jetson Thor, loses the suffix it asked for and comes back as sm_110.
The PTX ISA version follows the family: 5.0 for Maxwell and Pascal, 6.3 for Turing, 8.1 for Ampere and Ada, 8.5 for Hopper, 8.7 for sm_120, 8.8 for sm_100, sm_103 and sm_121, and 9.0 for Thor. The path that real programs use behaves the same way.
Built with mojo build --target-accelerator sm_90 --emit asm, a program that compiles the kernel through DeviceContext.compile_function writes a PTX file next to its assembly that says .target sm_90a, and the same build for sm_100 and sm_87 says sm_100a and sm_80.
NVIDIA’s assembler makes the consequences concrete, and its own help text states the rules. PTX for sm_XY compiles to any target at or above XY, suffixed or not. PTX for the family target sm_XYf compiles to members of the same family at or above it.
PTX for sm_XYa compiles to sm_XYa and nothing else. Every one of Mojo’s upgraded outputs therefore fails when handed to ptxas with the architecture that was requested, for example:
$ ptxas -arch sm_90 vadd_for_sm90.ptx
ptxas fatal : PTX with .target 'sm_90a' cannot be compiled for architecture 'sm_90'
$ ptxas -arch sm_100 vadd_for_sm90.ptx
ptxas fatal : Program with .target 'sm_90a' cannot be compiled to future architectureChange nothing but the directive back to .target sm_90 and the same body assembles for sm_100 and for sm_120. The same holds one generation later. Mojo’s output for a B200, sm_100a, is refused for the B300’s sm_103. Rewritten as the family target sm_100f, it assembles for sm_103 and is refused for sm_120, which belongs to a different family.
Of the twenty names Mojo accepts, twelve produce PTX that ptxas 13.4 will assemble for the architecture that was asked for. The three Maxwell and Pascal names fail earlier: ptxas 13.4 does not recognize sm_52, sm_60 or sm_61 as architectures at all.
Mojo lists them, and emits PTX 5.0 for them, but the current CUDA assembler cannot finish the job. Modular knows. A string inside the GPU runtime tells users of older hardware to point MODULAR_NVPTX_COMPILER_PATH at an external ptxas.
A target is a product, not an instruction set
The suffix is not produced by a transformation anywhere in the compiler. It comes from a table. In the shipped 1.1.0 the table lives in info.mojo, where the H100 record already says sm_90a.
On the main branch it has moved to Mojo/stdlib/std/_gpu/host/_builtin_targets.mojo, and the first thing that file does is map command-line names to products:
A10 ._with_cli_values[["sm_86"]],
A100 ._with_cli_values[["sm_80"]],
OrinNano ._with_cli_values[["sm_87"]],
L4 ._with_cli_values[["sm_89"]],
RTX4090 ._with_cli_values[[] ], # Could be ["sm_89"], but `sm_89` resolves to L4.
B100 ._with_cli_values[[] ], # Could be ["sm_100", "sm_100a"], but both resolve to B200.
B200 ._with_cli_values[["sm_100", "sm_100a"]],
H100 ._with_cli_values[["sm_90" , "sm_90a" ]],Each product carries a compilation target written as a kgen attribute, and the H100’s is where the letter comes from:
comptime _h100_target = CompilationTarget[
_mlir_value=__mlir_attr[
`#kgen.target<triple = "nvptx64-nvidia-cuda", `,
`stdlib_plugin = "cuda", `,
`arch = "sm_90a", `,
`features = "+ptx85,+sm_90a", `,
`tune_cpu = "sm_90a", `,
...So sm_90 is not an instruction set to Mojo. It is a lookup key that resolves to “an H100”, and an H100 is compiled as sm_90a with PTX 8.5, 132 streaming multiprocessors in its record.
We parsed the whole table: 44 device records and 43 target records, because the B100 and the B200 share one. sm_89 resolves to an L4, so an RTX 4090 cannot be selected by name. sm_86 resolves to an A10.
On main, the lookup itself is a compile-time Mojo program that walks a type list of target collections, and the list has two members: BuiltinTargets and one called ADDITIONAL_TARGETS. The device registry is pluggable at compile time, which we come back to at the end.
Whether the suffix is a bug depends on where the compiler is when the kernel gets built. MAX carries it at runtime. libmax.so contains the pipeline we timed above: the pass names LowerLIT, ElaborateGenerators, InlineParametric, VerifyParameters, CheckLifetimes and LowerKGENToLLVM are all in it, next to 255 strings that mention NVPTX and 466 that mention AMDGPU. When MAX compiles a kernel for the GPU in the machine, sm_90a unlocks every instruction the H100 has and costs nothing.
An executable made with mojo build is a different object. We built one for sm_90. It is 44,512 bytes, links only libKGENCompilerRTShared.so, libAsyncRTMojoBindings.so and the C library, and contains no pass names at all. What it does contain is one 1,354-byte PTX module that says .target sm_90a. NVIDIA’s assembler builds that module for sm_90a and refuses it for sm_100, sm_103 and sm_120 as a future architecture.
Change the one directive to sm_90 and it builds for all three. So the kernels of an H100 executable cannot be carried to Blackwell by NVIDIA’s compiler, where a plain target would have let the driver compile them forward, and nothing in the build said so.
Triton makes the same choice as Mojo: in our Triton piece, its sm_arch_from_capability appends the suffix to every capability from 9.0 upward, with a live TODO next to it. NVIDIA’s own Tile IR makes the opposite one.
Its tileiras backend accepts eleven targets and none with the suffix, and buys forward compatibility by excluding the instructions that would need it. Three serious projects, two good answers. Mojo and Triton recompile for the exact chip. cuTile narrows what a kernel can say.
The family targets sit between them, and the NVIDIA assembler bundled in max-core knows sm_100f, sm_103f, sm_110f, sm_120f and sm_121f. Mojo’s table asks for none of them.
So the cost of the choice depends on where portability has to live. In MAX it lives in the source plus a runtime that carries a full compiler inside the 162.6 MB libmax.so, with NVIDIA’s assembler beside it in the wheel.
In a mojo build executable it ends at build time, and the executable is exactly as portable as a CUDA binary compiled only for sm_90a.
Where the table is wrong
A device-centric design is only as good as the device table, and a hand-maintained table of 44 products accumulates small errors. We found four. They matter less individually than as a pattern.
The first explains Jetson Orin. The OrinNano record on main says arch = "sm_87", yet 1.1.0 emitted sm_80. The emitted LLVM IR already carried "target-cpu"="sm_80" with +ptx81,+sm_80, so the substitution happens in Mojo’s own code, before LLVM sees the kernel. The answer is in the 1.1.0 tag, where the table had a different structure.
Mojo/stdlib/std/_gpu/host/info.mojo maps "sm_87" to OrinNano correctly, then a method named target() turns a device into its compilation target through a chain of name comparisons that ends like this:
if self.name == "A100":
return _get_a100_target()
...
if self.name == "Jetson Thor":
return _get_jetson_thor_target()
...
if self.name == "":
return _get_empty_target()
return _get_a100_target()The chain has 43 name comparisons, and every product has a branch except one. There is no "Orin Nano" case, so the Orin falls through to the default, which is the A100.
Code built this way still runs on an Orin, because Ampere binaries are compatible within the family, so this is not a crash. It is a silent substitution, and the failure mode it represents is the one to worry about: an unrecognized device becoming an A100 instead of becoming an error.
Main has since restructured the table so that each device is declared together with its target record, as TargetAccelerator[GPUInfo, target], which removes this failure mode by construction. The shipped 1.1.0 still has it.
The second is an RTX 3090 record, unreachable from the command line because sm_86 resolves to the A10, whose features read +ptx63,+sm_86. PTX 6.3 predates sm_86. Compiling our kernel against that record, the way the standard library itself would, stops LLVM outright:
LLVM ERROR: PTX version 6.3 does not support target 'sm_86'. Minimum required
PTX version is 7.1. Either remove the PTX version to use the default, or increase
it to at least 7.1.Like the GTX 1080 Ti, GTX 1060, GTX 970 and Tesla P100 records, its name is NVIDIA’s full marketing string, "NVIDIA GeForce RTX 3090", which suggests these entries exist to match physical cards by name at runtime. We could not test whether a real 3090 selects it. The entry is identical at the 1.1.0 tag and on main.
The third is cosmetic: the RTX 4090 and 4090 mobile records compile for sm_89 but set tune_cpu to sm_90a, Hopper’s tuning, in both versions. In practice the NVPTX backend does little architecture-specific scheduling, and ptxas does the real work, so we expect no measurable effect.
The fourth is Jetson Thor, whose record asks for sm_110 without the suffix. That one does have consequences, and they show up with the tensor cores.
Where the tensor cores are
MAX exposes a warp-level matrix multiply-accumulate as one generic function, mma(d, a, b, c) in max.gpu.compute.mma, taking four SIMD values. Its body is a compile-time dispatch on the target: _mma_nvidia if the target is an NVIDIA GPU, _mma_amd if it is AMD, and _mma_apple only if the target is an Apple M5 with the metal4_0 feature.
We compiled a one-call kernel against every family and read what came out.
def k(o: Pointer[Float32, MutAnyOrigin]):
var a = SIMD[DType.bfloat16, 8](1)
var b = SIMD[DType.bfloat16, 4](1)
var c = SIMD[DType.float32, 4](0)
var d = SIMD[DType.float32, 4](0)
mma(d, a, b, c)
o[unsafe_offset=0] = d[0] + d[3]With those operand widths, sm_80, sm_90a, sm_100a and sm_120a all produce the same PTX instruction, mma.sync.aligned.m16n8k16.row.col.f32.bf16.bf16.f32. On Turing, sm_75, which has no bfloat16 tensor cores, compilation does not fail with a Mojo diagnostic. LLVM aborts: Cannot select: intrinsic %llvm.nvvm.mma.m16n8k16.row.col.bf16.
Change the fragments to four lanes each and the AMD targets compile: gfx90a emits v_mfma_f32_16x16x16bf16_1k, while gfx942 and gfx950 emit v_mfma_f32_16x16x16_bf16. RDNA3 rejects four-lane fragments inside its own implementation file and accepts sixteen-lane bfloat16 inputs with eight accumulators, emitting v_wmma_f32_16x16x16_bf16.
The M5 under Metal 4 wants eight lanes of fp16 in and eight of fp32 out, and emits the AIR intrinsic air.simdgroup_matrix_16x16x16_multiply_accumulate. Hand it four lanes and a constraint explains that Apple MMA requires eight-element fragments.
On an M4, a clean compile-time error says the target does not support the operation at all.
So one function name really does reach five distinct matrix instructions in four families of matrix hardware. It does not do so from one piece of code.
The operand widths are the vendor’s register fragment layout, and they differ on every family: eight, four and four lanes on NVIDIA for this shape, four, four and four on CDNA, sixteen, sixteen and eight on RDNA3, eight, eight and eight on Apple.
The abstraction that actually hides this is the layout package. Its TensorCore type is parameterized on output type, input type, an MMA shape and a transpose flag, and computes the fragment shapes for the target at compile time.
This is the precise sense in which Mojo is portable: the portability is in parameters evaluated at compile time, not in the operation. It is a good design. It also means that the thing a kernel author writes against is layout, which ships with MAX, not with Mojo.
The Hopper warpgroup multiply shows the other edge. wgmma_async compiled for sm_90a emits wgmma.mma_async.sync.aligned.m64n64k16.f32.bf16.bf16, and NVIDIA’s assembler accepts it. Compiled for sm_100a, sm_120a or sm_80, it produces no Mojo diagnostic. LLVM aborts on %llvm.nvvm.wgmma.commit_group.sync.aligned. The function’s compile-time asserts check shapes and scale factors, and we found none that checks the target.
It’s a third, independent confirmation of what we measured in our piece on Blackwell’s tensor memory, that wgmma exists on no Blackwell target at all. It also shows where Mojo’s knowledge of the hardware ends and LLVM’s begins.
The tcgen05 guard
Blackwell’s datacenter tensor cores are programmed through the tcgen05 family and a fifth memory space, tensor memory, which has to be allocated before use. MAX guards these instructions.
The file’s ten public tcgen05 functions make eleven calls to one constraint, identical at the 1.1.0 tag and on main:
# max/mojo/max/gpu/compute/arch/tcgen05.mojo
def check_blackwell_constraint():
comptime assert _has_blackwell_tcgen05(), (
"The tcgen05 instructions are only applicable on nVidia Blackwell"
" (sm_100a, sm_101a) hardware."
)
# Mojo/stdlib/std/sys/info.mojo
comptime _SM_101X_ARCHS: List[StaticString] = ["sm_101", "sm_101a"]
def _has_blackwell_tcgen05() -> Bool:
return _has_nvidia_gpu_any[
_SM_100X_ARCHS + _SM_101X_ARCHS + _SM_103X_ARCHS
]()
def _has_nvidia_gpu_any[archs: List[StaticString]]() -> Bool:
comptime if not has_nvidia_gpu_accelerator():
return False
comptime for arch in archs:
comptime if arch.removeprefix("sm_") in _accelerator_arch():
return True
return FalseTwo details in that code decide everything. The first is the choice of _has_nvidia_gpu_any. It asks whether the build has an NVIDIA accelerator configured, from the --target-accelerator flag or the detected GPU, and compares against that accelerator’s name. It does not ask what the kernel is being compiled for.
The same file has a sibling, _is_nvidia_gpu_any, that checks the compilation target. The second detail is the list: sm_101 was the name of Jetson Thor’s architecture before CUDA 13 renamed it sm_110, and Mojo’s own device table uses the new name. NVIDIA’s current assembler no longer accepts sm_101a at all.
We compiled a kernel that allocates tensor memory with every combination that matters, then gave each result to ptxas 13.4:
Kernel targetBuild acceleratorMojo 1.1ptxas 13.4sm_100anone (no flag, no GPU)refused by the guardnot reachedsm_100asm_100aemits tcgen05.alloc.cta_group::1.sync.aligned.shared::cta.b32acceptedsm_103asm_103aemits the sameacceptedsm_100asm_90arefused by the guardnot reachedsm_90asm_100aemits tcgen05 into sm_90a PTXrejected: tcgen05.alloc needs PTX ISA 8.6sm_110a (Thor)sm_110arefused by the guardaccepts tcgen05 on sm_110a and sm_110f, refuses it on sm_110sm_120a, sm_121asame as the targetrefused by the guardrefuses tcgen05 on both
On a GPU-less build machine, a common kind of CI runner, public tcgen05 code does not compile for a B200 unless someone passes --target-accelerator sm_100a. On a machine configured for an H100, a correct B200 kernel is refused. On a machine configured for a B200, the guard lets tcgen05 into Hopper PTX that NVIDIA’s assembler then rejects.
It’s possible that reading the build host is intentional, a way of instantiating Blackwell paths only when building for a Blackwell system, and that MAX’s own launch path always configures the accelerator to match.
The measured behaviour stands either way: in a language whose whole premise is compiling for hardware you are not running on, a hardware guard that consults the build host is the wrong primitive.
Jetson Thor gets both problems at once. Its architecture is missing from the guard’s list, which still says sm_101, and its device record compiles for sm_110 without the suffix, where ptxas refuses tcgen05 even when it is emitted. NVIDIA’s assembler accepts the identical instruction on sm_110a and sm_110f. As shipped, Mojo 1.1 cannot put Thor’s tensor memory to use, and fixing either blocker alone would not be enough.
The source explains why the Hopper and Blackwell failures looked so different: wgmma aborted inside LLVM, while tcgen05 sailed through LLVM into Hopper PTX.
They reach the instruction set by different roads. The warpgroup commit reaches LLVM as an NVVM intrinsic, the one named in LLVM’s error, and LLVM’s instruction selector checks intrinsics against the target. tcgen05.mojo contains 17 inline-assembly strings and no LLVM intrinsic calls: tcgen05.alloc is written out as text, and LLVM passes inline assembly through without reading it.
For tcgen05, the Mojo guard is the only check between the source and NVIDIA’s assembler. It is not an isolated choice. Across the max.gpu package we count 128 inline-assembly call sites against 76 LLVM-intrinsic call sites, alongside 296 compile-time asserts; across the kernel sources, 55 against 85, alongside 4,413 asserts.
Step back and there are three places where a portable stack can keep its knowledge of the hardware. The first is compile-time constraints written in the language, and when Mojo uses them the diagnostics are good: the M4’s clean refusal, the M5’s fragment-size message, the intent of the tcgen05 guard.
The second is LLVM’s instruction selector, which knows what each target can encode but can only say so by aborting the process. The third is the vendor’s assembler, which knows everything and speaks last. Inline assembly is a road around the middle layer: it goes from a Mojo string to the vendor’s assembler with nothing in between except whatever constraint the library author remembered to write.
A portable language is only as good as the first layer. In Mojo 1.1, a good part of the hardware knowledge still lives in the second and third.
The last mile belongs to the vendor
On AMD, Mojo owns the path to machine code. The AMDGPU backend runs in process, the assembly it prints is the final GCN or RDNA instruction stream, and the runtime loads the result through hipModuleLoadData from libamdhip64.so.
We found no reference to AMD’s comgr library in the runtime. On Apple, Mojo stops earlier and deliberately: its output for an M-series target is LLVM IR with the air64-apple-macosx triple and Metal feature flags, handed to Apple’s toolchain to finish.
On NVIDIA, Mojo stops at PTX, and the rest is NVIDIA’s. The GPU runtime, libMGPRT.so, and libmax.so both reference cuModuleLoadDataEx from libcuda.so.1, the driver’s own just-in-time path, and the entry points of NVIDIA’s PTX compiler library, nvPTXCompilerCreate and nvPTXCompilerCompile. The library they look for, named in the runtime’s own strings, is the 37.7 MB libNVPTX.so in max-core, and despite the name it has nothing to do with LLVM’s NVPTX backend.
It exports 45,501 dynamic symbols, and 45,474 of them are prefixed libnvptxcompiler_static_. The other 27 are NVIDIA’s public API: fifteen nvPTXCompiler functions, eleven nvLinker functions for linking device code, and one JIT entry point. Its version string reads Cuda compilation tools, release 13.3, V13.3.33, and it depends on nothing but the C and C++ runtimes.
We’re talking about ptxas as a library, NVIDIA’s static PTX compiler repackaged as a shared object and redistributed inside the MAX wheel. The runtime’s error strings name the two ways around it: MODULAR_NVPTX_LIBRARY_PATH for a different copy of the library, and MODULAR_NVPTX_COMPILER_PATH for an external ptxas binary, recommended for older drivers and older hardware.
This is the part of the stack we called the load-bearing moat in our piece on NVIDIA’s compiler: since Kepler and Volta, the dependency interlocks for fixed-latency instructions live in control bits that ptxas writes, so the assembler is correctness-critical and not merely an optimizer.
Mojo replaces CUDA’s source language and a growing share of its library layer, and on NVIDIA it is finished by the same closed assembler as a CUDA C++ kernel. That’s not a criticism of Modular, to be fair. It is the most direct measurement we have of how far a language can reach on NVIDIA hardware, and the answer is up to PTX, one step short of the instructions the chip executes.
We measured nothing about speed, so the best evidence on what that reach buys comes from outside. Godoy and colleagues at Oak Ridge National Laboratory, in a paper at the SC’25 workshops, ported four scientific workloads to Mojo, a seven-point stencil, BabelStream, miniBUDE and Hartree-Fock, and compared them with CUDA on an H100 and HIP on an MI300A.
They found Mojo competitive with the vendor baselines on the memory-bound kernels, and gaps on AMD for atomic operations and for fast-math compute-bound kernels. The study predates 1.0 by almost a year, and it fits the picture here: the language reaches each vendor’s backend, and what remains is how well each backend is driven.
The vendor libraries never left
MAX is usually described, including by us in July, as a stack that needs no vendor libraries: its own kernels, generated by its own compiler. max-core contains five precompiled binding packages, _cublas, _cudnn, _cufft, _rocblas and _miopen, 1.25 MB together.
Seven source files in MAX’s linear-algebra kernels reference vendor BLAS, including the Hopper and Blackwell matmul dispatchers. The Hopper dispatcher’s own comment says that on a miss, an unsupported configuration or no tuning for the shape, it returns a miss so the caller can fall back to vendor BLAS.
A helper named _vendor_blas_fallback_disabled turns that fallback off, and its documentation says it returns true only when the fallback has been disabled globally or a benchmark asks specifically for the Mojo kernel. The fallback is on unless someone turns it off.
The tuning is where the work is. The matmul tree contains 188 Hopper and 148 Blackwell tuning-config entries, most of them in lists keyed by exact problem shapes. One list is headed as the GEMM shapes of Llama 405B in FP8, with entries such as M=64, N=16384, K=2048.
That is the same structure we found in AMD’s hipBLASLt, where 16,177 kernels cover 2.9 million shape mappings: the moat is not the code generator but the campaign of measurements that decides which generated kernel to use for which shape.
The phrase “no vendor library” describes the tuned path. The default dependency graph still contains cuBLAS. We repeated Modular’s framing uncritically in our July piece on MLIR, and this is the correction.
How portable the kernels are
The language reaches four families of matrix hardware. The question left is how much of MAX’s kernel library does. We classified every Mojo file under max/kernels/src, 544,039 lines, by its path.
Files that live in a directory or carry a name for one vendor or architecture, such as sm90, sm100, amd_structured or apple, hold 156,621 NVIDIA-specific lines (28.8%), 93,998 AMD-specific lines (17.3%) and 11,903 Apple-specific lines (2.2%).
The remaining 281,517 lines (51.7%) sit on generic paths, and those still branch on the target at compile time. At least 48.3% of the kernel library is written for one vendor or one architecture.
That is not a failure of the language. It is what fast kernels look like, and CUTLASS keeps separate kernel families per architecture generation for the same reason.
What Mojo changes is that the NVIDIA, AMD and Apple slices are written in one language, against one layout algebra and one set of compile-time abstractions, and share the half of the library that is generic. It also gives the ModCon claim a yardstick.
If adding an accelerator really takes ten times less effort than it used to, the next vendor’s slice should look more like Apple’s 11,903 lines than AMD’s 93,998. That is our inference, and it becomes checkable the day that slice appears.
What Qualcomm bought, as far as the binary can tell
The price is public. Modular raised a $250 million round in September 2025 at a $1.6 billion valuation, bringing total capital to $380 million, according to RuntimeWire.
Qualcomm announced the acquisition on 24 June 2026 in an all-stock deal of about $3.9 billion and closed it on 29 July, with Chris Lattner named executive vice president of advanced AI software and platforms.
That is 10.3 times the capital raised and 2.4 times the last private valuation, the best outcome in the census of GPU software exits we assembled in August. Then came Mojo 1.0 on 11 August, the open compiler on 18 August, and Mojo 1.1 on 17 September, which also opened the compiler to outside contributions.
At ModCon, according to the NAND Research’s account, Modular announced platform support for AWS Trainium, Google TPUs and Qualcomm’s Cloud AI100 and Dragonfly accelerators, said that integrating a new accelerator now takes more than ten times less engineering effort than before, and announced an alliance program for hardware vendors, model providers and clouds.
None of those accelerators appears in what we measured. Mojo 1.1.0 lists NVIDIA, AMD and Apple GPUs and nothing else. No source file in the open repository mentions Qualcomm; the word occurs only inside four test-data files, two GGUF model files and two tokenizer vocabularies.
Hexagon appears in a comment in a C++ support file, in a test fixture, and in a documentation page whose sample target listing includes it next to amdgcn. The shipped binary lists neither: its host backends are AArch64, RISC-V and x86.
The mechanism for adding them is visible, though. On main, the target lookup walks two collections of devices, the builtin one and ADDITIONAL_TARGETS. A vendor can supply its own collection of device records at compile time without editing the table everyone else uses.
Our reading, which is inference and not measurement, is that the announced backends live outside the open tree, at least for now, or reach customers in a form we cannot see, and that the plumbing for plugging them in is already public.
The claim of ten times less effort is testable in exactly one way: the day a Qualcomm target appears in --print-supported-accelerators or in the open tree, the size of the change that added it will say how much effort it took.
What this says about the deal is the more interesting question. Nearly everything we counted is open: the compiler, the standard library and MAX’s kernels, all 523 kernel source files carrying the same Apache 2.0 with LLVM exceptions header as the compiler.
What stays closed is the runtime that compiles and runs models, the 162.6 MB libmax.so and the GPU runtime beside it, shipped under the MAX license, along with whatever target collections live outside the tree.
Opening the compiler commoditizes the frontend, much as NVIDIA let its own become a commodity years ago, building NVVM on LLVM and contributing the NVPTX backend upstream.
Qualcomm is not buying a way to make kernels portable; roughly half of them are written per architecture anyway. It is buying the team, the closed runtime and the plumbing that turns its own accelerators into one more target collection, next to a kernel library that was tuned on other people’s hardware first.
The repository’s root already carries governance for that future: an AI tool policy that asks contributors to label assisted work with an Assisted-by trailer, to keep pull requests small because AI lowers the cost of generating code and not of reviewing it, and to keep a human in the loop, alongside AGENTS.md and CLAUDE.md files written for coding agents.
If you build on Mojo today
Build one executable per GPU architecture. A mojo build for sm_90 embeds sm_90a PTX that NVIDIA’s compiler will not build for Blackwell, and the Hopper outputs and nearly all the Blackwell ones are locked the same way. The family targets such as sm_100f would carry a B200 build to a B300, and Mojo does not ask for them.
The path that follows the hardware is MAX’s, which compiles on the machine the model loads on. On a CI runner without a GPU, pass --target-accelerator with the architecture your kernels target, or public tcgen05 code will not compile, and never build Hopper kernels on a runner configured for Blackwell.
On Jetson Orin, 1.1.0 compiles as sm_80; on Jetson Thor it cannot use tensor memory; on Maxwell and Pascal, point MODULAR_NVPTX_COMPILER_PATH at a ptxas old enough to know them.
If you need a guarantee that no vendor library is in the path, build with -D MODULAR_DISABLE_VENDOR_FALLBACK=true, the define the matmul dispatch reads to turn off the cuBLAS and rocBLAS fallback.
Executables built from a pip install carry an absolute RUNPATH into that Python environment. And don’t contort code to avoid Mojo’s 64-bit Int in index math: on our vector kernel, Int32 index math saved 2 of 20 PTX instructions and 4 of 11 64-bit operations, ptxas allocated 10 registers either way, and the cubin shrank by 128 bytes.
Three of the problems above, the Orin fallthrough, the RTX 3090 PTX version and the two Thor blockers, are small patches. Since 17 September the compiler accepts outside contributions, which may be the most useful thing the open-sourcing did.
Where I might be wrong
The tcgen05 guard may read the build host on purpose. If MAX’s own launch path always sets the build accelerator to the device a kernel targets, the mismatch we produced can only happen when someone compiles by hand the way we did.
The consequence we measured, tcgen05 inside Hopper PTX, stands, but its practical reach may be small.
We compiled through _compile_code and, to check target selection, through DeviceContext.compile_function, and both picked the same targets. We never launched a kernel. Launching applies rules we never exercised, such as the requirement since 1.0 that kernel arguments be fixed-width types rather than Int. PTX that assembles is not code that is correct or fast.
Nor did we load a mojo build executable on a Blackwell GPU: the refusal of its embedded PTX is the verdict of NVIDIA’s assembler, which the runtime would also have to get past, not something we watched happen.
The kernel census classifies files by path. It is a lower bound on architecture-specific code, and a path that names a vendor can still hold helpers other vendors use.
Our timings come from a one-core container and small programs. Absolute numbers will differ on real machines. We expect the ratios to hold, the LLVM share especially, but the specialization curve is a result about small bodies only.
A PyPI metadata field is not a license text. We did not read the MAX license in full, and we are not lawyers.
The RTX 3090 record may be unreachable in practice. And backends absent from the open tree may exist and ship through channels we cannot see, such as MAX containers or partner builds. The claim is “not in the open tree”, not “does not exist”.
Five dated predictions
First: by 31 March 2027, a stable Mojo release removes at least one of the two Jetson Thor blockers, either by adding sm_110 names to the tcgen05 guard or by compiling Thor as sm_110a or sm_110f.
Second: through 31 December 2027, the default target records for the H100, B200 and B300 stay architecture-specific, with the a suffix. MAX’s runtime compiler has no reason to give those instructions up, and we expect mojo build users to be told to build per architecture rather than see the default change.
Third: by 30 June 2027, a Qualcomm accelerator appears in --print-supported-accelerators of a public Mojo release or in the open repository.
Fourth: through 31 December 2027, every shipped configuration of Mojo and MAX that runs on NVIDIA hardware still finishes code with NVIDIA’s closed PTX compiler, whether bundled, invoked through the driver’s JIT, or as an external ptxas.
Fifth: fn does not return in any 1.x release.
Confidence dossier
Every claim that carries weight in the piece, graded. M means we measured it and verify.py re-checks it; A means we read it in a primary source, usually the code at a named revision; B means it was reported or claimed and we did not reproduce it; C is our inference. 42 rows are M, 1 reuse our own earlier measurements, 15 are A, 4 are B and 6 are C.
Reproduce it
Everything above ran in a one-core x86-64 container with 3 GB of memory and no GPU, on 21 and 22 September 2026. The probe scripts and the harness are published with this piece as mojo-compile-probe, MIT licensed. verify.py holds 187 checks and exits non-zero at the first number that no longer matches.
pip install mojo==1.1.0 max-core==26.6.0 # Mojo 1.1.0 (8189361e), MAX 26.6.0
pip download nvidia-cuda-nvcc==13.4.92 --no-deps # ptxas V13.4.92, unzip and use bin/ptxas
git clone --depth 1 https://github.com/modular/modular.git # measured at 26cfe94
git -C modular fetch --depth 1 origin tag mojo/v1.1.0 # c6fa49f, what shipped
python3 m2_target_sweep.py # 30 targets, emitted .target and PTX version
python3 m3_ptxas_roundtrip.py # every NVIDIA output through ptxas, forward-compat tests
python3 m4_tensor_cores.py # mma, wgmma, tcgen05 per target and build flag
python3 m5_language.py # the language probes
python3 m6_compile_time.py # specialization scaling and timing census
mojo build aot.mojo --target-accelerator sm_90 -o aot_sm_90 # then extract the PTX and run ptxas
mojo build m7_launch_path.mojo --target-accelerator sm_90 --emit asm # launch-path check
python3 verify.py --live # re-checks every number quoted; exits non-zero on failureCorrections before publication
Writing the checks before the prose caught these, in the order they happened. We list them because the method only works if its failures are visible.
While probing imports, a shell redirection error made it look as if wildcard imports of missing modules were silently accepted. They are not; the compiler reports them correctly. Caught by re-running the probe with separate output files.
The family targets sm_100f through sm_121f were first attributed to Mojo’s LLVM NVPTX backend. The strings came from libNVPTX.so, which is NVIDIA’s nvptxcompiler, not LLVM.
The harness caught a conflation in the symbol count: libNVPTX.so exports 45,501 dynamic symbols, of which 45,474 carry the libnvptxcompiler_static_ prefix.
A chart subtitle said MAX’s code was five times the compiler. The measured ratio for MAX’s Mojo code is 3.3.
The Jetson Orin downgrade was first diagnosed against main, whose table is structured differently. The diagnosis in the text is against the v1.1.0 tag that actually shipped.
A draft said the RTX 3090 record was the only one named with a full marketing string. The GTX 1080 Ti, GTX 1060, GTX 970 and Tesla P100 records are too.
The first specialization-scaling run used one measurement per point and called the curve flat. Three runs per point show a small but real slope, about 0.56 ms per specialization; the text and Figure 4 use the repeated data.
Version 2 said the suffix costs nothing because Mojo compiles at load time and never ships PTX. That holds for MAX, whose runtime library carries the compiler. It is false for mojo build executables, which embed build-time PTX and no compiler; the targets section now separates the two.
Tuning-config counts included the struct definitions: the call sites are 188 and 148, not 189 and 149.
The pass count mixed 2 passes from the shared Support library into the Mojo compiler’s 45.
Hexagon also appears in a documentation page’s sample target list, not only in a comment and a test fixture.
Version 2 explained the compile floor as the standard library’s cost without measuring it; the empty-program census now does.
Final read-through: the opening said sm_90a unlocks the tensor memory accelerator. ptxas 13.4 accepts TMA copies on plain sm_90; the suffix is what wgmma and setmaxnreg need, and the text now names those.
Final read-through: version 3 said 1.1 deleted two keywords. The release notes list three keywords (fn, alias, __comptime_assert) and two syntax forms.
Final read-through: version 3 contrasted an open language with a kernel library it implied was not. All 523 kernel source files carry the same Apache 2.0 header; the closed part is the runtime, and the Qualcomm section now says so.
Final read-through: a sentence left over from version 2 still called the specialization curve flat.
Two framings from our July piece on MLIR are corrected in the text: KGEN is a dialect inside Mojo rather than a layer above it, and MAX’s matmul path keeps a vendor BLAS fallback.
Sources
Measurements, analysis and conclusions are ours. Source excerpts from modular/modular are Apache 2.0 with LLVM exceptions; compiler and assembler messages are output we generated. Charts and harness: mojo-compile-probe.