DEV Community: Masahiro Morino

How Does Rust Code Run on a GPU? Tracing cuda-oxide and PTX

Masahiro Morino — Sun, 01 Nov 2026 11:40:58 +0000

I used cuda-oxide to run a vector addition kernel written in Rust on a GPU.

Pass in two arrays and get back their element-wise sums. The calculation is simple, but I wanted to understand what form that Rust function takes before it reaches the GPU.

Among the files generated during the build was some code called PTX. This time, I will read it to see how Rust's thread indexing and array accesses are represented.

What is cuda-oxide?

cuda-oxide is part of NVIDIA's CUDA Rust project. It provides a compiler and related tools for writing kernels for NVIDIA GPUs in Rust. See the official CUDA Rust repository.

The official book introduces it this way:

cuda-oxide is a Rust-to-CUDA compiler that lets you write (SIMT) GPU kernels in safe(ish), idiomatic Rust.

Source: The cuda-oxide Book.

In other words, you can write the actual function that runs on the GPU in idiomatic Rust and compile it. SIMT stands for Single Instruction, Multiple Threads. It is the programming model used in this example, where many threads execute the same kernel.

As of October 2026, CUDA Rust describes itself as an early-stage research project, with its features and APIs still under development. Here, I will focus on a small vector addition example and look at what the Rust code becomes. See the project status.

As someone who writes Rust, I find this a very intriguing project.

Adding arrays: who computes each element?

Here is the kernel we will inspect. A kernel is a function that serves as an entry point for execution on the GPU.

I verified the code and output on November 1, 2026, under the following conditions. I built the vecadd example from cuda-rust commit a0cc6cc0bc18722ec5cdfbea136c118a7c051977 in a source tree separate from my ongoing changes.

Item	Verified environment or setting
Host	Linux, x86-64
GPU	NVIDIA GeForce RTX 5060 Ti (Compute Capability 12.0)
NVIDIA driver	595.99.02
Rust toolchain	nightly-2026-08-28, rustc 1.100.0-nightly (e457a7b0d)
LLVM	Version 23.1.0 bundled with the toolchain
Build	Release, target sm_120a, ordinary PTX output with bounds checks enabled

use cuda_device::{DisjointSlice, cuda_module, kernel, thread};

#[cuda_module]
mod kernels {
    use super::*;

    #[kernel]
    pub fn vecadd(a: &[f32], b: &[f32], mut c: DisjointSlice<f32>) {
        let idx = thread::index_1d();
        let idx_raw = idx.get();
        if let Some(c_elem) = c.get_mut(idx) {
            *c_elem = a[idx_raw] + b[idx_raw];
        }
    }
}

This is based on the vecadd example included in NVIDIA's cuda-rust repository. The original code is published under the Apache-2.0 license. See the cuda-rust source.

The calculation is c[i] = a[i] + b[i]. However, there is no for loop inside the function to traverse the whole array.

On a CPU, we could compute the elements sequentially like this:

for i in 0..n {
    c[i] = a[i] + b[i];
}

In this GPU kernel, each thread works out the index of the element it is responsible for and computes that one element. Many threads execute the same kernel.

Thread responsible for element 0: c[0] = a[0] + b[0]
Thread responsible for element 1: c[1] = a[1] + b[1]
Thread responsible for element 2: c[2] = a[2] + b[2]
...

thread::index_1d() produces this index. Its return value, idx, is not just an integer: it has cuda-oxide's ThreadIndex type. idx.get() extracts the integer used to index the input arrays.

On the output side, DisjointSlice accepts this thread index and provides a mutable reference to the element assigned to the thread. If c.get_mut(idx) returns Some, the thread writes its result to that element. I will leave the details of these types aside for now and focus on which elements each thread reads and where it writes.

On the host, meaning the CPU side, the program transfers the input data to GPU buffers, specifies the thread configuration, launches the kernel, and retrieves the result. We need to distinguish the whole program, including that preparation, from the kernel running on the GPU.

PTX is a representation before GPU machine code

So what is the PTX generated from this kernel?

PTX stands for Parallel Thread Execution. It is a virtual instruction set for NVIDIA GPUs. It is not specific to Rust: compilers for languages such as CUDA C++ generate PTX too. See the NVIDIA PTX specification.

PTX contains instruction names and register-like names, but another conversion step remains before we have the machine code the actual GPU executes. You will also encounter the name SASS for the instructions of the physical GPU.

Here, I will use PTX to see how the Rust expression is broken down into operations. I will not treat it as output that already fixes the final machine-instruction count or physical-register allocation.

The vecadd.ptx file starts with these directives:

.version 8.7
.target sm_120a
.address_size 64

They specify the PTX version, target GPU architecture, and address width, respectively. PTX has version and target requirements, so the same file is not necessarily usable with every NVIDIA GPU or driver.

After these directives come the kernel entry point, .visible .entry vecadd(, its parameters, and register-variable declarations. Let's look at what happens in the body.

Where does cuda-oxide connect to rustc?

Before reading the PTX, let's establish the path from Rust source to this point.

cuda-oxide provides a compiler, libraries, and build tools for writing CUDA kernels in Rust. At its center is rustc-codegen-cuda, a custom code generation backend for rustc.

It uses rustc for tasks such as parsing Rust, type checking, and borrow checking. GPU code generation takes the resulting MIR and proceeds by transforming it into other representations. See the cuda-oxide architecture.

The ordinary PTX generation path we are following can be simplified like this:

Rust kernel
    | rustc performs parsing, type checking, borrow checking, and other work
    v
MIR
    | cuda-oxide transforms it using Pliron IR
    v
LLVM IR
    | LLVM's llc generates code with the NVPTX backend
    v
PTX

Pliron is a framework for building and transforming intermediate representations. cuda-oxide uses it to handle representations of Rust MIR and lower them to LLVM IR.

LLVM produces the PTX at the end of this path. cuda-oxide writes out LLVM IR and passes it to LLVM's code generation tool, llc. Here too, LLVM provides the infrastructure for generating code for a different target.

We can break down the roles as follows:

Name	Role in this path
rustc	Rust analysis, checking, and MIR processing
rustc-codegen-cuda and related crates	Extract GPU-side work and transform MIR into LLVM IR
LLVM's llc and NVPTX backend	Generate PTX from LLVM IR
cargo-oxide	Provide commands for building and related operations

Host-side code in the same file, such as main, goes to the ordinary LLVM backend and becomes CPU code. In this example, the host executable also contains the GPU's PTX.

Calling the body of a #[kernel] function as an ordinary Rust function on the CPU does not move execution onto the GPU. The host uses a launch API to invoke it as a GPU kernel.

Finding one line of Rust in the PTX

For this verification, I built and ran the example from the cuda-oxide workspace, specifying the example and target. --verbose also displays information such as the tools used to generate LLVM IR and PTX.

cargo oxide run vecadd --arch sm_120a --verbose

The final output was as follows. The program compared the results with values calculated on the CPU, and all 1024 elements matched.

✓ SUCCESS: All 1024 elements correct!

The same build generated crates/rustc-codegen-cuda/examples/vecadd/vecadd.ptx. The following excerpts come from that file, with whitespace adjusted for readability.

Working out the element index

First, let's find the instructions related to thread::index_1d(). Omitting an intervening parameter-load instruction gives us this sequence:

mov.u32      %r1, %ctaid.x;
mov.u32      %r2, %ntid.x;
mov.u32      %r3, %tid.x;
mul.wide.u32 %rd11, %r1, %r2;
cvt.u64.u32  %rd12, %r3;
add.s64      %rd13, %rd11, %rd12;

CUDA groups threads into blocks and arranges those blocks in a grid. In one dimension, the global element index is the block index multiplied by the number of threads per block, plus the thread's index within its block. See CUDA's thread hierarchy.

The first three mov instructions read the values used in that calculation.

Special register	Value read
`%ctaid.x`	Block index in the x direction
`%ntid.x`	Number of threads per block in the x direction
`%tid.x`	Thread index within the block in the x direction

For example, with 256 threads per block, block index 2, and thread index 7 within the block, the element index is 2 * 256 + 7 = 519.

In the PTX, mul.wide.u32 multiplies two 32-bit values and places the 64-bit result in %rd11. The next cvt converts the thread index within the block to 64 bits, and add.s64 adds the values. %rd13 now holds the element index.

What was a single function call in Rust has become reads of the block and thread coordinates, a multiplication, and an addition.

Names such as %r1 and %rd13 refer to virtual register variables in PTX. We do not need to interpret their numbers as physical GPU register numbers.

From an element index to a memory location

Next come the array accesses.

The index we just calculated is in %rd13, but the later instructions use %rd14. This instruction sits between them:

selp.b64 %rd14, %rd13, -1, %p7;

selp selects between two values according to a condition. If %p7 is true, it places the value of %rd13 in %rd14; otherwise, it selects -1 as an invalid value. In this output, %p7 represents the one-dimensional launch condition: the y and z dimensions of both the grid and the block must all be 1.

When that condition holds, %rd14 contains the element index calculated earlier. One instruction that uses it is:

shl.b64 %rd15, %rd14, 2;

shl is a left shift. Here, it shifts the element index left by two bits, multiplying it by four.

Each element of these arrays is an f32, which occupies four bytes. For element 519, the distance from the start of the array is 519 * 4 = 2076 bytes.

The PTX also adds that distance to the starting addresses. The relevant instructions are:

add.s64 %rd3,  %rd8, %rd15;
add.s64 %rd18, %rd1, %rd15;
add.s64 %rd19, %rd2, %rd15;

In this output, %rd1 points to the start of input a, %rd2 to the start of input b, and %rd8 to the start of output c. Adding the byte offset produces the addresses to read from and write to.

The compact expression a[idx_raw] hides the element size and the address calculation.

Read two values, add them, and write back

Now that we have the addresses, let's look at the calculation. In the actual output, input-address calculations alternate with loads.

add.s64       %rd18, %rd1, %rd15;
ld.global.b32 %r8, [%rd18];
add.s64       %rd19, %rd2, %rd15;
ld.global.b32 %r9, [%rd19];
add.f32       %r10, %r8, %r9;
st.global.b32 [%rd3], %r10;

ld.global.b32 loads 32 bits from global memory. Square brackets, as in [%rd18], indicate an access to the value at that address. The two inputs are loaded into %r8 and %r9.

Next, add.f32 performs a 32-bit floating-point addition, placing the sum in %r10. The .b32 loads handle 32-bit bit patterns, which the addition interprets as f32 values. There is no integer-to-floating-point conversion in between.

Finally, st.global.b32 writes the result to the output address in %rd3. Unlike a load, the bracketed destination appears on the left.

We can now see the work corresponding to this line of Rust:

*c_elem = a[idx_raw] + b[idx_raw];

The addition itself is one instruction, but the surrounding code computes the element index, calculates memory addresses, and reads and writes the values.

Bounds checks remain as instructions too

Around that calculation are instructions that test conditions and branch.

For example, before accessing input a, we find:

ld.param.b64 %rd4, [vecadd_param_1];
setp.ge.u64  %p13, %rd14, %rd4;
@%p13 bra   $L__BB0_5;

In this output, vecadd_param_1 is the length of input a. It is passed as a parameter separate from the starting address.

setp.ge.u64 tests whether the element index in %rd14 is greater than or equal to the length in %rd4 and places the result in the predicate variable %p13. The next instruction, @%p13 bra, branches to the named label if that condition is true.

What is at the destination?

$L__BB0_5:
    trap;

trap aborts execution and notifies the host of an error. In this output, an out-of-bounds access to input a takes this path. Input b has a similar check. The bounds checks associated with Rust's indexing are visible as PTX comparisons and branches. See the PTX trap instruction.

The output is handled differently: the Rust code writes only when c.get_mut(idx) returns Some. Skipping the calculation when the output index is out of bounds and trapping when an input index is out of bounds are two different paths.

The number of GPU threads can be slightly larger than the number of elements. For example, launching four blocks of 256 threads for 1000 elements gives 1024 threads. The extra 24 threads must not write past the end of the array. These index checks have a practical role.

What remains after PTX is generated?

We have read the PTX, but the journey to GPU execution is not finished.

On the path used here, where PTX is passed to the CUDA driver for loading, the driver's JIT compiler translates it into machine code for the target GPU. It can also cache the compiled result for reuse. See NVIDIA's explanation of JIT compilation.

The overall flow is:

Build time
Rust kernel -> MIR -> Pliron IR -> LLVM IR -> PTX
Rust host code -> CPU code
                      |
                      v
           Embedded in the host executable

Run time
The host program starts
    | Pass PTX to the CUDA driver and load the module
    v
Machine code for the target GPU (using JIT compilation or its cache)
    | Launch the kernel from the host
    v
GPU threads compute and write results to memory
    |
    v
The host waits for completion and retrieves the results

There are also paths that build a GPU binary from PTX ahead of time. Here, I am following the path that loads PTX.

Running Rust code on a GPU involves more than sending the Rust source to it. There is checking in rustc, transformation in cuda-oxide, PTX generation in LLVM, and then GPU-specific compilation and launch work.

What we found in the PTX was thread-index acquisition, element-address calculation, two loads, one addition, and a store. That short array expression is becoming readable as the work each GPU thread performs.

What Is LLVM?

Masahiro Morino — Sun, 25 Oct 2026 06:06:39 +0000

When you look into how Rust is compiled, the name LLVM comes up.

Here, I will compare it with Clang and GCC, familiar from C and C++, to clarify LLVM's role and see what rustc delegates to LLVM.

LLVM is infrastructure for building compilers

LLVM is not the name of a programming language. It is a collection of reusable components and technologies for building compilers and related tools. See the official LLVM overview.

The name LLVM can refer to the entire project or to its optimization and code generation infrastructure. When we say that Rust compilation "passes code to LLVM," we mainly mean the latter.

That includes tasks such as transforming a calculation into a more efficient form and selecting instructions for ARM64 or x86-64. Language implementations can share this infrastructure instead of building those parts from scratch.

LLVM IR, on the other hand, is a language used during compilation. IR stands for Intermediate Representation. It helps to distinguish LLVM, the infrastructure, from LLVM IR, the representation of code that it processes.

Name	What it refers to
LLVM	Compiler infrastructure and tools; depending on context, its optimization and code generation infrastructure in particular
LLVM IR	The intermediate representation LLVM uses for optimization and code generation
rustc	The Rust compiler, which normally uses LLVM for code generation
Clang	A compiler for C, C++, and related languages that uses LLVM
GCC	A collection of compilers for multiple languages, with its own optimization and code generation infrastructure

We will return to Clang and GCC later. First, I want to separate the infrastructure from the representation of code used inside it.

Compilation is more than replacing one instruction with another

"Turning source code into machine code" might suggest replacing each line of Rust with CPU instructions in sequence.

In practice, several stages sit between the two.

The compiler parses the source, resolves what names refer to, and checks that types fit together. For Rust, it also needs to check ownership and borrowing rules.

It then transforms the program through different representations, removes unnecessary computations, and selects instructions suitable for the target CPU. Eventually, it determines registers and instruction sequences and generates machine code.

The part that handles a language's syntax and meaning is often called the frontend; the part that analyzes and optimizes intermediate representations, the middle end; and the part that generates target-specific code, the backend. Actual implementations do not always have neat boundaries between those three. See the rustc overview and LLVM code generator documentation.

LLVM provides shared components for tasks including optimization and code generation. Using LLVM does not automatically give a compiler an understanding of Rust's syntax or ownership rules.

What does rustc delegate to LLVM?

When using the usual LLVM backend, rustc analyzes and checks Rust source, processes intermediate representations such as MIR, and generates LLVM IR.

MIR is an intermediate representation on the Rust side. It is used for borrow checking and Rust-side optimization too. The division of work is not "leave every optimization to LLVM." See Rust's MIR.

LLVM then optimizes the LLVM IR and generates code for the target CPU. The resulting object files are passed to a linker. See rustc code generation.

Rust source code
    ↓ rustc parses, type-checks, and more
Rust-side intermediate representations (including MIR)
    ↓ Borrow checking, optimization, preparation for code generation, and more
LLVM IR
    ↓ LLVM optimizes and generates code for the target CPU
Object files
    ↓ Linker
Executables and other linked outputs

This is a simplified view of the division of work. Intermediate representations and stages are omitted.

"Passing code to LLVM" does not normally mean that rustc writes a text file ending in .ll and asks another command to read it. rustc's LLVM backend uses LLVM's API to construct and process LLVM IR in memory. A .ll file is a human-readable rendering of that IR. See LLVM IR representations and rustc's LLVM IR generation.

The compiler we use for Rust is rustc. LLVM is shared infrastructure incorporated into part of its work.

What does addition look like after it reaches LLVM?

Let's compile a small function to make this more concrete. The outputs below were produced with rustc 1.96.0 on an Apple Silicon Mac.

Save this as add.rs:

pub fn add(a: u32, b: u32) -> u32 {
    a.wrapping_add(b)
}

The function adds two u32 values. wrapping_add wraps the result around if it does not fit in a u32. I use it here because I want to observe the addition itself without including panic handling for overflow. See u32::wrapping_add.

The following command outputs LLVM IR, assembly, and an object file:

rustc add.rs --crate-name llvm_intro --crate-type lib --edition 2024 \
  -C opt-level=2 -C link-dead-code --emit=llvm-ir,asm,obj

--crate-type lib treats the source as a library without a main. -C opt-level=2 enables optimization so we can inspect the shortened computation. Since this example has no caller, I also use -C link-dead-code to keep the function's code from being removed. We stop before linking and request the three output formats explicitly. See rustc output options and code generation options.

Opening llvm_intro.ll, I found this definition of add. The file header and attribute definitions outside the function are omitted.

define hidden noundef i32 @_ZN10llvm_intro3add17he381aecb720e834bE(i32 noundef %a, i32 noundef %b) unnamed_addr #0 {
start:
  %_0 = add i32 %b, %a
  ret i32 %_0
}

Set the long function name and attributes aside for a moment. Start with these two lines:

%_0 = add i32 %b, %a
ret i32 %_0

i32 is a 32-bit integer type, and %a and %b are the argument values. add adds them, names the result %_0, and ret returns it.

The Rust u32 type and method-call syntax are not preserved as they appeared in the source. The operation is now expressed as "add two 32-bit values and return the result." LLVM IR's i32 type itself does not distinguish signed from unsigned integers; operations and other details express those distinctions where needed. This wrapping addition is represented by add i32 without additional flags. See LLVM IR integer types and the add instruction.

An LLVM IR add is not a CPU add instruction

Next, open the generated llvm_intro.s. The function's instructions are:

add w0, w1, w0
ret

This is ARM64 assembly. w0 and w1 are 32-bit CPU registers, and add adds the two values and puts the result in w0. Here, the computation uses the registers holding the arguments, and the return value is placed in w0 too.

Wait, LLVM IR also had an add. They are not the same thing, though.

LLVM IR values such as %a and %_0 are not CPU register names. LLVM reads the relationships between these values and lowers them into instructions and registers for the target CPU. In this case, that produced ARM64's add w0, w1, w0.

We can inspect the machine code in llvm_intro.o too. This command uses llvm-objdump from Xcode Command Line Tools:

xcrun llvm-objdump --disassemble --demangle llvm_intro.o

Here are the instruction lines from its output:

0: 0b000020     add w0, w1, w0
4: d65f03c0     ret

From left to right, each line shows the instruction's location, its machine-code encoding, and a readable representation of the instruction. Our LLVM IR addition has become an instruction the CPU can execute.

LLVM IR was not the machine code itself. It represented the computation at an earlier stage.

LLVM IR is sometimes described as CPU-independent, but that does not make it universally portable code you can move unchanged to any environment. This .ll file includes a target triple identifying the target and a target datalayout describing data layout. Target-specific details such as calling conventions and pointer sizes also matter. See LLVM IR data layout.

Why doesn't rustc implement everything itself?

It is possible to implement machine-code generation specifically for Rust. But handling the instructions and register constraints of each supported CPU, while continuously improving optimization, requires substantial implementation and maintenance work.

Using LLVM lets rustc translate Rust's operations into LLVM IR and share the optimization and CPU-specific code generation that follow.

Clang uses LLVM too. Clang handles the meaning of C and C++, then uses infrastructure shared with Rust from LLVM IR onward. See Clang's official site.

Figure 1: Language-specific processing and the parts shared through LLVM. The arrows show the use of common infrastructure, not a merger of separate programs' IR into one program.

For example, when LLVM's optimizations improve, compilers for multiple languages can incorporate that improvement. Sharing CPU-specific code generation also reduces the burden of maintaining a separate implementation for each language.

Of course, adding support for a CPU in LLVM does not mean Rust automatically supports everything about it. Rust-side target configuration, standard library support, and other work are still needed. The benefit is the large amount of infrastructure that can be shared. See why rustc uses LLVM.

How do LLVM and GCC differ?

Both LLVM and GCC have mechanisms for sharing optimization and CPU-specific code generation across multiple languages. They provide separate implementations for the task of turning source code into machine code.

LLVM provides optimization and code generation functionality as libraries. Clang and rustc translate their respective languages into LLVM IR and use that infrastructure.

GCC stands for GNU Compiler Collection. It includes language-specific processing, intermediate representations, optimization, and CPU-specific code generation. GCC also shares processing across multiple languages, but its implementation is separate from LLVM's. GIMPLE, for example, is one of the intermediate representations it uses for optimization. See GCC's official site and GCC's GIMPLE documentation.

What we are comparing	LLVM	GCC
Language-specific processing	Clang, rustc, and other compilers handle their respective languages	GCC's language-specific frontends handle it
Intermediate representations used for optimization	LLVM IR	GIMPLE and others
CPU-specific code generation	Uses LLVM's backends	Uses GCC's backends

The distinction is not that only LLVM offers reusable components while GCC can only be used as a monolithic command. GCC also has libgccjit, a library that lets external programs use code generation and related functionality. See the libgccjit documentation.

When writing C or C++, choosing between Clang and GCC means choosing a compiler as a whole, from language processing through machine-code generation. Choosing Clang means using LLVM's infrastructure for optimization and code generation. See Clang.

For Rust, rustc is the compiler handling the language. It does not ask Clang to read Rust. rustc and Clang each handle their own languages while using LLVM's infrastructure.

Can Rust use GCC?

Can Rust code generation use GCC too?

The usual rustc configuration uses LLVM, but there are efforts to use GCC. It is useful to distinguish replacing rustc's code generation backend from implementing a Rust frontend within GCC. See rustc code generation backends.

The first approach is rustc_codegen_gcc. It keeps rustc's parsing, type checking, borrow checking, and other Rust-side functionality, while using GCC instead of LLVM for code generation. This backend uses libgccjit, mentioned above. Despite JIT in its name, the project uses it for ahead-of-time compilation: generating machine code before execution. See rustc_codegen_gcc.

The second is gccrs. Instead of replacing rustc's backend, this project implements a Rust frontend within GCC and connects it to GCC's optimization and code generation. See gccrs.

Configuration	The side that parses and checks Rust code	Optimization and code generation infrastructure
Usual rustc	rustc	LLVM
rustc with rustc_codegen_gcc	rustc	GCC through libgccjit
gccrs	Rust frontend within GCC	GCC

This table shows the division of work. As discussed earlier, rustc also performs optimizations on its side.

One goal of rustc_codegen_gcc is to make Rust available on platforms that LLVM does not support. Using GCC opens another path for target support and optimization. Which one is faster, however, needs to be evaluated for a particular program and target environment.

As of October 2026, these efforts remain under development. That does not mean they have the same language-feature coverage or setup requirements as the usual rustc configuration. For example, the rustc_codegen_gcc README says it requires a patched libgccjit. Here, the focus is on their different structures rather than installation instructions. See rustc_codegen_gcc requirements and gccrs development.

Rather than "Rust doesn't use GCC," the situation is that Rust normally uses LLVM, while alternative configurations using GCC are also being developed. Rust the language and LLVM the machine-code generation infrastructure are separate things.

LLVM's role in Rust compilation

The LLVM project as a whole also includes tools beyond Clang, such as the linker LLD and the debugger LLDB. That is why looking up LLVM brings up tools beyond machine-code generation. See LLVM subprojects.

To understand Rust compilation, it helps to distinguish rustc, which handles Rust's rules; LLVM IR, which represents the intermediate computation; and LLVM, which performs further optimization and code generation.

We may type just one command, rustc, but inside it, Rust-specific mechanisms work together with infrastructure shared with other languages.

What Happens Before Rust Calls Your main Function?

Masahiro Morino — Sun, 18 Oct 2026 09:00:31 +0000

When we write a Rust program, fn main() is where our code begins.

Does the CPU jump straight into that function? Who puts the executable into memory, and where does execution actually begin?

Let's inspect a small Rust program and follow what runs before the main we wrote. Looking inside the executable, it turns out there is more than one main.

This experiment uses an Apple Silicon Mac and an ordinary executable that uses the standard library, std. My environment is rustc 1.96.0 on macOS 26.6.2, with tools from Xcode Command Line Tools. Entry mechanisms, internal function names, and addresses can change across operating systems and Rust versions.

An executable needs some preparation before it can run

An executable on disk contains machine code and data. When the CPU executes instructions, though, it uses addresses in the process's memory.

At startup, the OS reads information from the executable and maps its code and data into the process's virtual address space. This does not mean copying the entire file into memory before anything can run. The system can establish mappings and load individual pages as they are needed.

For this macOS executable, the dynamic linker, dyld, also runs. It makes shared libraries available, performs the required reference resolution and initialization, and passes control to the executable's entry point. See the Mach-O format definitions and dyld startup implementation.

The dynamic linker works when the program starts. The build-time linker combines object files into an executable; the dynamic linker loads the shared libraries that executable uses and connects references to functions and other symbols.

Once that preparation is done, do we finally reach our own main?

Figure 1: The path we will follow in this executable. Helper functions and details such as library initialization are omitted. The dyld and Rust startup routines are themselves code already running on the CPU.

Build an executable with an ordinary main

No special attributes or unsafe are needed. Save this as main.rs:

fn main() {
    println!("Hello from Rust main");
}

Compile and run it:

rustc main.rs --edition 2024 \
  -C opt-level=0 -C debuginfo=2 --emit=link,llvm-ir
./main

Hello from Rust main

--emit=link,llvm-ir produces both the executable, main, and the LLVM IR, main.ll. I disable optimization for our code to make the calls easier to observe, and include debug information so a debugger can follow functions and source lines. This does not disable optimization in the already-built standard library. See rustc output options and code generation options.

Two main functions?

Disassemble the executable. --demangle makes Rust's encoded symbol names easier to read.

xcrun llvm-objdump --disassemble --demangle main

The output is long. For now, here are just two function headings:

00000001000009f4 <main::main::h2dd6e7c25b480450>:
0000000100000a2c <_main>:

Wait. Two mains? And at different addresses?

In main::main, the first main is our crate's name. The second is the function we defined. A hash is appended to distinguish symbols. Inside this function, I found the code that prints our message.

The other function, _main, is an entry wrapper generated by the compiler for this target. In Mach-O symbol output, its name has a leading underscore.

The fn main() in our source and the externally called entry function are separate functions. We did not write the entry wrapper ourselves.

Which one does the executable identify as its entry point?

Names alone do not establish which function is the entry point. Let's inspect the executable's metadata too.

otool -l main

Mach-O has load commands that describe information needed for memory mapping and startup. These are relevant excerpts from my output:

  segname __TEXT
   vmaddr 0x0000000100000000
  fileoff 0

          cmd LC_LOAD_DYLINKER
      cmdsize 32
         name /usr/lib/dyld (offset 12)

       cmd LC_MAIN
   cmdsize 24
  entryoff 2604
 stacksize 0

LC_LOAD_DYLINKER identifies the dynamic linker, while LC_MAIN records the entry location. Here, entryoff gives the entry's byte offset in __TEXT. See Mach-O's entry_point_command.

2604 in hexadecimal is 0xa2c. In this executable, __TEXT starts at file offset zero and has a preferred virtual address of 0x100000000, giving:

0x100000000 + 0xa2c = 0x100000a2c

That matches _main, not the 0x1000009f4 address of our main::main.

These are addresses in the layout recorded in the file. ASLR can shift the actual load addresses at runtime. Also, LC_MAIN does not mean "the first instruction the CPU executes at startup." It identifies this executable's entry point, reached after preparation by dyld and other startup code.

The entry wrapper does not call our main directly

Run the same disassembly command and extract the section starting with <_main>:.

xcrun llvm-objdump --disassemble --demangle main

0000000100000a2c <_main>:
100000a2c: a9bf7bfd    stp  x29, x30, [sp, #-0x10]!
100000a30: 910003fd    mov  x29, sp
100000a34: aa0103e2    mov  x2, x1
100000a38: aa0003e8    mov  x8, x0
100000a3c: 93407d01    sxtw x1, w8
100000a40: 90000000    adrp x0, 0x100000000 <dyld_stub_binder+0x100000000>
100000a44: 9127d000    add  x0, x0, #0x9f4
100000a48: 52800003    mov  w3, #0x0
100000a4c: 97ffff8f    bl   0x100000888 <std::rt::lang_start::hf66e9cda02e72e37>
100000a50: a8c17bfd    ldp  x29, x30, [sp], #0x10
100000a54: d65f03c0    ret

Each instruction line contains its address, the machine-code encoding, and a readable representation of the instruction. For example:

100000a34:  aa0103e2    mov x2, x1
Address     Encoding    Instruction and operands

mov x2, x1 copies the value in register x1 into x2. Registers are locations inside the CPU that hold values and addresses temporarily.

We do not need to decode every instruction. Start with bl, the instruction that calls a function. Under the calling convention used here, the first argument is passed in register x0.

The call goes to std::rt::lang_start, rather than our main. Before that, adrp and add construct 0x1000009f4 in x0. That was the address of our main::main. Focus on the address calculation here, rather than the symbolic annotation printed beside adrp.

The wrapper passes the location of our main as an argument to the standard library's startup routine. In effect: "Start this Rust program using this function."

We can check the same relationship in LLVM IR. Here is the entry wrapper from main.ll:

define i32 @main(i32 %0, ptr %1) unnamed_addr #4 {
top:
  %2 = sext i32 %0 to i64
; call std::rt::lang_start
  %3 = call i64 @_ZN3std2rt10lang_start17hf66e9cda02e72e37E(ptr @_ZN4main4main17h2dd6e7c25b480450E, i64 %2, ptr %1, i8 0)
  %4 = trunc i64 %3 to i32
  ret i32 %4
}

Inside the parentheses on the line calling lang_start, the first argument is:

ptr @_ZN4main4main17h2dd6e7c25b480450E

ptr is the pointer type, and @_ZN4main4main… refers to our main::main function. This argument is a pointer to that function.

Replacing the long names with readable descriptions gives the following argument order. This is explanatory notation, not literal LLVM IR:

lang_start(pointer to our main, argc, argv, setting)
           ↑ first argument

The remaining arguments are the command-line argument count, argc; argv, which points to the argument strings; and a setting related to signal handling.

Our own main has no declared parameters. The entry wrapper still handles information passed in from the operating system.

What does lang_start prepare?

The entry wrapper passed our main to std::rt::lang_start instead of calling it directly. What does lang_start do before calling our function? Let's look at the standard library implementation.

In library/std/src/rt.rs shipped with my Rust toolchain, lang_start is defined as follows. Its attributes are omitted.

fn lang_start<T: crate::process::Termination + 'static>(
    main: fn() -> T,
    argc: isize,
    argv: *const *const u8,
    sigpipe: u8,
) -> isize {
    lang_start_internal(
        &move || crate::sys::backtrace::__rust_begin_short_backtrace(main).report().to_i32(),
        argc,
        argv,
        sigpipe,
    )
}

The first parameter, main: fn() -> T, is the function pointer received from the entry wrapper. lang_start wraps the work of calling that function and converting its result into an exit code in a closure, then passes that closure to lang_start_internal. The &move || ... part creates the closure and passes a reference to it.

Next, here is an excerpt from the body of lang_start_internal. Comments mark omitted code; this excerpt is not independently compilable.

// Excerpt from the body of lang_start_internal
panic::catch_unwind(move || {
    unsafe { init(argc, argv, sigpipe) };

    let ret_code = panic::catch_unwind(main).unwrap_or_else(move |payload| {
        // Panic handling omitted
        101
    });
    let ret_code = ret_code as isize;

    cleanup();
    // Coordination of concurrent thread exits omitted

    ret_code
})
.unwrap_or_else(handle_rt_panic)

It first calls init, then runs the supplied closure through panic::catch_unwind(main). That leads to our own main. Afterward, it calls cleanup and returns the exit code, ret_code.

Leaving the detailed error handling aside, the sequence is:

lang_start_internal
    ↓
Initialize Rust runtime state
    ↓
Invoke the supplied closure, reaching our main
    ↓
Handle the return value as an exit code
    ↓
Clean up stdout and other runtime state
    ↓
Return the exit code to the entry wrapper

Initialization records the main thread's ID and proceeds to platform-specific preparation. In the Unix implementation used here, that includes checking the standard I/O file descriptors, configuring SIGPIPE behavior, and preparing stack overflow handling. On macOS, it also sets the main thread's name. See std startup and Unix initialization.

SIGPIPE is a signal relevant to situations such as writing to a pipe whose reader has gone away. Startup handles settings we rarely think about when using println!.

Rust does not prepare everything from scratch, though. The Apple-specific command-line argument implementation uses information already prepared by dyld. Rust's preparation builds on work done by the OS and dynamic linker. See the Apple argument implementation.

When main returns normally, execution proceeds to cleanup. The call is also wrapped in catch_unwind to handle panic unwinding. With a configuration that aborts on panic, however, the process terminates without unwinding.

Asking whether Rust "has a runtime" can be ambiguous. What we can see here is that this ordinary executable using std contains startup routines implemented in Rust.

Stop inside main and look at the path we took

We can also check this while the program runs. Use LLDB to place a breakpoint in our own main:

lldb ./main

Inside LLDB, run:

breakpoint set -n main::main
run
bt

bt walks the current call stack: the current function, its caller, that caller's caller, and so on. My program stopped at the println! line. See the LLDB tutorial.

Here are the relevant frames, with long arguments, hashes, and line numbers removed. The frame numbers are unchanged.

frame #0:  main::main
frame #1:  core::ops::function::FnOnce::call_once
frame #2:  std::sys::backtrace::__rust_begin_short_backtrace
frame #3:  std::rt::lang_start::{closure}
           (frames #4 through #11 omitted)
frame #12: std::rt::lang_start_internal
frame #13: std::rt::lang_start
frame #14: main
frame #15: dyld`start

The current location is at the top, so read from bottom to top to follow the calls that brought us here.

From dyld, through the entry main, lang_start, and lang_start_internal, we finally reach our own main::main. Along the way are a closure and helper functions used to invoke it.

Now "code runs before our main" is something we can see in an actual call stack.

Returning from main is not the final step either

Our main has no explicit return type, so it returns (). lang_start uses the Termination trait's report method to convert that result into an exit status. For (), the status indicates success. See Termination.

The startup code does more than invoke our function: it translates a Rust return value into a result for the program as a whole. After standard library cleanup, the entry wrapper returns an exit code, leading to the system's process termination handling.

There are other ways to terminate, such as process::exit and aborting, which do not return through this entire path. We followed the normal completion of this small program.

Even for a three-line program, the OS and dyld prepare it, and Rust's standard library performs initialization, before the main we wrote is called.

Why Rust Object Files Still Need a Linker: Symbols and Relocations

Masahiro Morino — Sat, 10 Oct 2026 10:59:01 +0000

When you compile Rust code, object files are produced along the way. They have the extension .o and contain machine code and data. The linker combines them to build an executable or another output.

But something bothers me here. If the machine code is already there, shouldn't the CPU be able to execute it? Why is linking still necessary?

Let's compile a pair of small functions separately and look at a state where the machine code exists, but the destination of a function call has not been determined yet. Symbols and relocations will help us read that state.

I tried this on an Apple Silicon Mac. The object format is Mach-O, and the instructions are ARM64. On Linux or Windows, the format and inspection commands differ. The addresses and instruction output below were obtained on October 10, 2026.

Can we generate machine code without the callee's implementation?

In everyday Rust, we define functions with fn and use Rust's usual mechanisms to call functions in other crates. We do not need to write unsafe or extern "C" just to use the linker.

For this experiment, I want to build the caller and callee as separate .o files and compare the state without the callee's implementation with the state after linking. That calls for a slightly unusual setup. extern "C" gives both sides the same calling convention, and #[unsafe(no_mangle)] makes the symbol names easy to follow. The "C" specifies a calling convention; it does not mean we are using C source code. Both implementations are written in Rust.

With an external function declared this way, however, the compiler alone cannot verify that the declaration matches the actual definition. That is why unsafe appears in this declaration and call. It is part of the setup chosen to observe the linker, rather than a requirement for ordinary Rust function calls.

First, let's write a function that simply calls answer. Save this as caller.rs:

#![no_std]

unsafe extern "C" {
    fn answer() -> u32;
}

#[unsafe(no_mangle)]
pub extern "C" fn call_answer() -> u32 {
    unsafe { answer() }
}

We have the body of call_answer, but not the body of answer. All we have for answer is a declaration saying that it takes no arguments and returns a u32.

Since this looks a little different from everyday Rust, let's go through the settings used for the experiment:

The unsafe extern "C" block declares an external function. We specify its type and calling convention, but this declaration alone does not let the compiler verify that they match the actual definition.
extern "C" selects the C ABI calling convention. Here, it gives the caller and the definition the same convention. We will write the implementation in Rust later.
#[unsafe(no_mangle)] disables the usual symbol name mangling so that we can follow the function names easily. It is an unsafe attribute because the names can collide with other symbols.
#![no_std] prevents this small library from automatically using the standard library, std. The main function we write later will use std.

For more about these settings, see the Rust Reference sections on External blocks, The no_mangle attribute, and The no_std attribute.

The following command produces an object file before linking:

rustc caller.rs --crate-type lib --edition 2024 \
  -C opt-level=0 -C panic=abort --emit=obj -o caller.o

--emit=obj requests a native object file. --crate-type lib tells rustc to compile code without a main function as a library. Here, we are taking the .o requested by --emit=obj, rather than an .rlib. We also specify -C panic=abort, so unwinding on panic is outside this experiment. Command-line Arguments: --emit

Even without the body of answer, caller.o was produced.

file caller.o

caller.o: Mach-O 64-bit object arm64

So this part works. But what happened to the call to answer?

Object files contain more than machine code

An object file is more than a sequence of machine instructions. Alongside code and data, it carries information needed to connect them later.

We will focus on three things:

Contents	What we will examine
Code section	What the machine code for `call_answer` looks like
Symbol table	Which names this file defines, and which it needs from elsewhere
Relocation information	Where references must be adjusted, and which targets they refer to

A section is a region that holds a particular kind of content, such as code or data. In my caller.o, the instructions were in the __text section of __TEXT. Here, text means code, rather than written text.

Figure 1: The three kinds of information we examine in caller.o. This diagram does not show their actual positions or sizes in the file. Headers and other information are omitted.

Mach-O also has headers and load commands that let tools locate this information. We will read the parts relevant to the function call rather than decode the entire file format. On macOS, nm inspects the symbol table, and otool inspects sections and other contents. Building Mach-O Files — Apple

The symbol table records the missing implementation

First, let's use nm to see the externally visible symbols:

nm -g caller.o

                 U _answer
0000000000000000 T _call_answer

A symbol is information used to refer to a function or another entity by name. In this Mach-O output, the source names answer and call_answer appear with a leading underscore.

T denotes a definition in the code section; U denotes a symbol that is undefined in this file. In other words, call_answer is here, while the definition of answer must come from elsewhere. llvm-nm: symbol types

At this stage, U is not a failure. The object file records the intention to call something defined elsewhere.

The 0000000000000000 to the left of _call_answer does not mean the function will execute at address zero, either. In this .o, it is at the start of the code section. Its placement in the final executable is still to be determined.

So the name is still there. What about the machine code?

There is a call instruction. But where does it go?

Let's disassemble the code with otool. Disassembly turns machine code bytes back into a human-readable instruction listing.

otool -tvV caller.o

caller.o:
(__TEXT,__text) section
_call_answer:
0000000000000000 stp x29, x30, [sp, #-0x10]!
0000000000000004 mov x29, sp
0000000000000008 bl  0x8
000000000000000c ldp x29, x30, [sp], #0x10
0000000000000010 ret

The instruction to focus on is bl at 0000000000000008. ARM64 uses this instruction for function calls, and it also preserves information needed to return. Around it, we can see register saves and restores associated with the call. Arm: BL — Branch with Link

The 0000000000000008 on the left is the address where the bl instruction sits. The bl 0x8 on the right shows its branch destination as 0x8.

Wait. The instruction is at 0x8, and its destination is also 0x8? Read on its own, this looks like an instruction that branches to itself.

This bl instruction encodes the distance from the instruction to its destination. That distance is called a relative displacement. In this output before linking, it is still a placeholder value of zero. So 0x8 + 0 = 0x8, and the disassembler shows a branch to the same location. During linking, the linker writes the correct distance to answer into the instruction.

What we actually want to call is answer. To see that intention, we need to read the instruction together with the relocation information.

Relocation information says: adjust this reference to answer

Use otool -rv to display the relocations:

otool -rv caller.o

caller.o:
Relocation information (__TEXT,__text) 1 entries
address  pcrel length extern type    scattered symbolnum/value
00000008 True  long   True   BR26    False     _answer

Let's match this row to the instruction we just saw:

Field	Meaning in this example
`address: 00000008`	Apply the relocation 8 bytes from the start of the `__text` section
`pcrel: True`	Treat the reference as relative to the instruction's position
`type: BR26`	A relocation that adjusts the displacement of an ARM64 branch instruction
`_answer`	The symbol used to determine the target

Here, address does not mean byte 8 from the start of the entire file. It is an offset within the code section. The file also contains headers and other information, so section offsets and file offsets must be read separately.

BR26 corresponds to Apple's ARM64_RELOC_BRANCH26, which handles the 26-bit displacement in B and BL instructions. Different instruction types need different adjustments, so a relocation needs both the referenced symbol and a relocation type. Apple: arm64 relocation types

The symbol table tells us which names exist. The relocation information tells us where to adjust references to those names.

That lets us produce the call instruction and the instructions for completing it later, even before we know where answer will be placed.

The word “relocation” might sound as though we are simply moving a function somewhere else. What we are examining here is information for adjusting a reference encoded in an instruction to match the final placement.

Put the body of answer in another .o

Now let's write the callee. Save this as provider.rs:

#![no_std]

#[unsafe(no_mangle)]
pub extern "C" fn answer() -> u32 {
    42
}

Compile this into a separate object file as well:

rustc provider.rs --crate-type lib --edition 2024 \
  -C opt-level=0 -C panic=abort --emit=obj -o provider.o

nm -g provider.o

0000000000000000 T _answer

This time, _answer has a T. The definition that caller.o needs now exists in provider.o.

But this address is also zero. call_answer is at zero, and answer is at zero. Are they in the same place?

No: they are at the starts of their own sections in separate object files. During linking, they are placed in the output, and their references are connected using that placement.

Figure 2: A conceptual view of resolving the reference to answer. Other linker inputs, such as main and the standard library, are omitted.

Matching a name's reference to its definition is symbol resolution. Adjusting the necessary references according to the layout is the relocation processing we are looking at here.

Putting the files next to each other does not connect them. We need to establish which code calls which destination, and where that destination is.

Leave out the implementation, and linking fails

Let's add an entry point so we can run the program. Save this as main.rs:

unsafe extern "C" {
    fn call_answer() -> u32;
}

fn main() {
    println!("{}", unsafe { call_answer() });
}

This main uses std. First, deliberately add only caller.o to the build:

rustc main.rs --edition 2024 -C opt-level=0 -C panic=abort \
  -C link-arg=caller.o -o demo-missing

This time, linking failed. Here is the actual error, with only the long linker command and the note about its arguments omitted:

error: linking with `cc` failed: exit status: 1
  |
  … (linker command and argument note omitted)
  = note: Undefined symbols for architecture arm64:
            "_answer", referenced from:
                _call_answer in caller.o
          ld: symbol(s) not found for architecture arm64
          clang: error: linker command failed with exit code 1 (use -v to see invocation)

error: aborting due to 1 previous error

Undefined symbols means that a symbol's definition could not be found. The missing name is _answer, and the error identifies _call_answer in caller.o as the code referencing it.

We have call_answer, but cannot find the answer it calls. A reference that could remain unresolved when producing the .o cannot be resolved when building this executable, so linking fails.

Now let's also provide provider.o:

rustc main.rs --edition 2024 -C opt-level=0 -C panic=abort \
  -C link-arg=caller.o -C link-arg=provider.o -o demo

./demo

It works.

-C link-arg adds an argument to the linker invocation. Here, we use it to add the two .o files. rustc also supplies inputs needed for main, the standard library, and so on. These two files are not the executable's only ingredients. Codegen Options: link-arg

The instruction really changed after linking

Finally, let's compare the machine code before and after linking. I used the LLVM disassembler available on my machine:

xcrun llvm-objdump --disassemble caller.o
xcrun llvm-objdump --disassemble demo

Below are just the lines containing the call to answer. From left to right, they show the address, the machine instruction's value, and its human-readable form. Tabs and column spacing have been normalized.

Before linking, in caller.o:

8: 94000000  bl 0x8 <ltmp0+0x8>

After linking, in demo:

100030a0c: 94000003  bl 0x100030a18 <_answer>

94000000 became 94000003. This is more than a function name being added to the display: the value encoded in the call instruction itself changed. llvm-objdump

In my output, call_answer was placed at 0x100030a04, and answer at 0x100030a18. Given that layout, the call instruction at 0x100030a0c now branches to answer. We can see the relocation applied to the actual machine code.

The “3” encoded in the instruction does not mean 3 bytes. The 26-bit displacement in bl (imm26) is a signed value in units of 4 bytes. Here, it means 3 × 4 = 12 bytes ahead: 0x100030a0c + 12 = 0x100030a18, the address of answer. Arm instruction specification

The addresses shown here are virtual addresses recorded in the executable. They are distinct both from byte offsets within the file and from addresses necessarily used unchanged on every run. ASLR can shift the layout at load time, and references to shared libraries also involve runtime mechanisms. Here, we have focused on a call between two functions included in the same executable. Executing Mach-O Files — Apple, Apple Platform Security: ASLR

Why do we still need a linker when the machine code exists?

Our caller.o contained the machine code for call_answer. But answer remained an undefined symbol, accompanied by relocation information for adjusting the call instruction to reference it.

The compiler producing caller.o can generate an instruction saying “call answer here.” At that point, however, we have not supplied provider.o. We do not know where the body of answer will sit in the final executable, so we cannot determine the distance to it yet.

Instead, we leave behind the target's name and the location of the instruction that needs adjustment. Once the necessary inputs are available, the linker does the following:

Match references to definitions. It connects the answer requested by caller.o to the definition in provider.o. When we left out the implementation, resolution failed and produced an error.
Determine the layout of code and data. It places sections and other contents from separate .o files into the output executable. Positions local to each object file become part of the final layout.
Adjust references to match that layout. It reads the relocation information and, in our example, writes the displacement to answer into the bl instruction. This is why 94000000 became 94000003.

Simply concatenating the files would leave the zero in the call instruction unchanged. Even if we included the function's body, the call would not be connected unless we adjusted the reference to reach it. Producing machine code and combining that code into a working program are separate jobs. Building Mach-O Files — Apple

Then why not compile everything together from the start? With a small example, that is tempting. But real programs use the standard library and other libraries alongside our own code. Compiling pieces separately and combining them later allows existing machine code to be reused. There is a qualification in Rust: generic functions, for example, may be monomorphized for concrete types and have machine code generated in the consuming crate. Not everything can be supplied solely as precompiled machine code. When producing the final executable, references between the generated pieces still need to be resolved. Monomorphization

The extern "C" in this experiment made that boundary easy to observe. Even without this syntax in everyday Rust, separately generated code still has to be connected. rustc can divide a single crate into multiple Codegen Units; with the LLVM backend, the object files produced from those units are linked. This normally happens as part of a build through rustc or Cargo, so we rarely need to operate the linker directly. Code generation

There was still work left between “the machine code exists” and “we have something executable.” The undefined symbol and relocation information we saw were how that work was handed to the next stage.

What Code Does Rust Pass to LLVM? Generics and Codegen Units

Masahiro Morino — Sat, 03 Oct 2026 05:05:26 +0000

In the previous article, I generated MIR from Rust code and tried reading it. An if expression became a set of blocks, and we could see where values went and which block execution would enter next.

But being able to read MIR doesn't mean we've reached machine code yet. LLVM apparently comes into the picture further along. What does rustc actually pass to it, and how?

For example, if we call one generic function with both u32 and f64, does LLVM receive one function? Or two? Even if both use x + x, integer addition and floating-point addition involve different operations, don't they?

As a follow-up to Reading Rust's MIR, this article explores which code rustc passes to LLVM, and how it groups that code.

The key ideas are monomorphization of generics and partitioning into Codegen Units (CGUs). We'll start with collecting the code we need, then look at dividing it into units, and finally at converting it to LLVM IR. I started out investigating linkers, and somehow I still haven't reached one.

This article is based on the Rust Compiler Development Guide as consulted on October 1, 2026. I rechecked the examples with Rust 1.96.0, the same version used in the previous article, and used that version's implementation when looking at CGU partitioning. Compiler internals aren't fixed by the Rust language specification.

Getting a feel for the path after MIR

rustc normally uses LLVM as its code-generation backend. On the Rust side, rustc handles syntax, types, borrowing, and related processing, then lowers the program to an intermediate representation called LLVM IR. LLVM receives that representation and performs optimization and machine-code generation for the target environment. Here, we'll follow the path that uses the LLVM backend. Code generation — Rust Compiler Development Guide

Once MIR is ready, the part of the process we're interested in looks like this:

Figure 1: A conceptual view from MIR to object files. Details such as optimization across units through LTO are omitted.

Let's pause here. MIR, LLVM IR, CGU. With these names lined up, they all start to look like similar intermediate representations.

Name	What it represents
MIR	A representation of Rust operations for analysis and transformation
LLVM IR	A representation of code that rustc passes to LLVM
CGU	A grouping that determines which code undergoes code generation together

MIR and LLVM IR are formats for representing code; a CGU is a unit that groups code-generation items. One concerns how code is represented, and the other concerns which code is processed together. They appear in sequence in the diagram, but they're describing different things.

rustc first collects the functions and other items it needs, then decides which CGU to put them in. It lowers MIR to LLVM IR for each grouping and passes the result to LLVM. Let's look at what the “collection” and “partitioning” in the middle of the diagram are for.

Collecting the concrete types that are needed

First comes collection: identifying the items needed for code generation.

In a generic function, a type parameter such as T doesn't simply become a machine-code type as written in the source. Rust generates code for the combinations of concrete types it needs. This is monomorphization.

For example, this twice function adds a value to itself:

fn twice<T: std::ops::Add<Output = T> + Copy>(x: T) -> T {
    x + x
}

fn main() {
    let _ = twice(3_u32);
    let _ = twice(1.5_f64);
}

Add<Output = T> requires the addition's result to have type T as well. Copy lets us use the same x on both sides of the addition. This example uses the function in two ways: with T = u32 and with T = f64.

Here, I'll call a function with concrete type arguments, such as twice::<u32>, an “instance.” We're talking about a function to generate code for, rather than a value or object stored in a variable.

Figure 2: A conceptual view focused on twice. Other code-generation items, such as main and the addition implementations, are omitted. The two columns do not imply that the instances go into separate CGUs.

It helps to distinguish collecting the required items from generating code for their concrete types.

The monomorphization collector first identifies the instances needed for code generation, such as twice::<u32> and twice::<f64>. This includes ordinary functions such as main and statics, too. The items are then placed in CGUs, and their type parameters are treated as concrete types during lowering from MIR to LLVM IR. Monomorphization, Lowering MIR to a Codegen IR

When I emitted LLVM IR for this code locally, I found two definitions of twice: one taking i32 and another taking double. The commands and environment are listed later, under “Generating LLVM IR locally.”

Use in Rust	Argument and return types in this LLVM IR output
`twice::<u32>`	Both `i32`
`twice::<f64>`	Both `double`

Wait, i32 for a Rust u32? In LLVM IR, i32 means a 32-bit integer type. It doesn't carry the same “signed” meaning as Rust's i32. Signedness is distinguished by the operations or their specifications, such as the kind of comparison being performed. Similar-looking names don't always mean the same thing. LLVM Language Reference: Integer Type

Counting “one function definition” in source code is therefore different from counting “how many concrete function instances are needed” during code generation. This is where writing one function with T connects to performing operations on particular types.

That doesn't mean two functions always survive in the final binary. This example doesn't use the results, so optimization may remove the calls altogether. With inlining and other transformations in the picture, we need to distinguish the collected items from the machine code that eventually remains.

CGUs group the code passed to LLVM

Once the required code-generation items have been collected, rustc partitions them into Codegen Units (CGUs). With the LLVM backend, each CGU gets an LLVM module: a grouping of LLVM IR. An LLVM module doesn't correspond one-to-one with a Rust mod.

Here's a conceptual example with two CGUs:

Figure 3: A conceptual view of processing two CGUs. LLVM's work in the two columns can proceed in parallel. LTO and duplication of code for inlining are omitted.

Dividing the work LLVM can handle allows compilation to use multiple cores. Code generation

Another consideration is incremental compilation. It reuses results from an earlier compilation where changes haven't affected them. For code-generation artifacts, rustc decides whether reuse is possible at the CGU level. If it can reuse a result, it can skip repeating code generation and optimization for that part. Ideally, we wouldn't regenerate every bit of code on every build. That also influences how CGUs are partitioned. CGU partitioning implementation and design notes, Rust 1.96.0

So should we just split everything into smaller and smaller units?

There are tradeoffs. When LLVM optimizes within one module, it can consider the functions and other items inside that module together. Partitioning affects both parallelism and the granularity of reuse, as well as opportunities for optimization across functions.

For example, inlining puts the body of a called function into its caller, which requires that body to be available. As we divide the work, we also have to consider how much code can be optimized together. A question about compilation speed has turned into a question about the speed of the generated program.

CGU partitioning	Potential benefit during compilation	What to watch for in the generated code
More CGUs	May allow more parallel LLVM work and finer-grained reuse	Partitioning can affect optimization opportunities and may reduce runtime performance
Fewer CGUs	Allows optimization over a wider scope within one module	Compilation may be slower, and better runtime performance isn't guaranteed

rustc's -C codegen-units=N sets an upper limit on the number of CGUs. The official documentation likewise explains that increasing the number may speed up compilation while producing slower code. Codegen Options: codegen-units

Optimization can also cross CGU boundaries. LTO (Link-Time Optimization) allows optimization across module boundaries. To understand runtime performance, we need to consider optimization levels and LTO settings alongside the CGU count. Codegen Options: lto

CGU boundaries don't simply follow source files

Looking at the diagrams, you might wonder whether rustc creates one CGU for each Rust source file. The relationship isn't that simple.

During incremental compilation, Rust 1.96.0's partitioning process separates non-generic code associated with a module from instances of generic functions. The generic instances can change as the combinations of types used change, even when the function bodies themselves haven't been edited.

For example, we could leave the definition of twice alone but add a use of twice::<u64> elsewhere. That would require an instance for the new type. Separating these changes from a module's non-generic code helps preserve parts that can be reused.

This is only part of the partitioning policy, though. The implementation also merges units to stay within the CGU count limit. We can't conclude that every module always produces exactly two final CGUs. CGU partitioning implementation, Rust 1.96.0

CGUs are groupings chosen with code generation and reuse in mind, so their organization can differ from how we organize the source.

What gets translated from MIR to LLVM IR?

Finally, let's look at the step toward LLVM IR in Figure 1.

The developer guide introduces rustc_codegen_ssa::mir::codegen_mir as the entry point for lowering MIR. The processing of MIR's elements is divided into modules such as these:

Module	What it handles
`block`	Basic blocks and terminators, including branches and function calls
`statement`	Statements such as assignments
`operand`	Values used in computations and calls
`place`	References to locations where values are stored
`rvalue`	Computations and other expressions producing a value to assign

In the previous article, we used choose to look at branches and their merge in MIR:

pub fn choose(flag: bool, x: u32) -> u32 {
    if flag { x } else { 0 }
}

For this function, LLVM IR needs to express which block execution enters depending on the condition, what goes into the return value, and where the function returns.

Let's emit choose.ll by adding --emit=llvm-ir to the same example. Omitting the function header with its name and arguments, the braces enclosing the function, and the comments beside the block labels, the body looked like this:

start:
  %_0 = alloca [4 x i8], align 4
  br i1 %flag, label %bb1, label %bb2

bb2:
  store i32 0, ptr %_0, align 4
  br label %bb3

bb1:
  store i32 %x, ptr %_0, align 4
  br label %bb3

bb3:
  %0 = load i32, ptr %_0, align 4
  ret i32 %0

There are more symbols now, but the branch from the previous article is still recognizable. Let's start with br, store, and ret.

LLVM IR line	What it does in this example
`br i1 %flag, label %bb1, label %bb2`	Enters `bb1` if `%flag` is true, or `bb2` if false
`store i32 %x, ptr %_0, align 4`	Writes `%x` to the memory pointed to by `%_0`
`br label %bb3`	Enters `bb3` unconditionally
`%0 = load i32, ptr %_0, align 4`	Reads a value from that memory and names it `%0`
`ret i32 %0`	Returns the value that was read

In start, alloca allocates space for the return value. The two branches write to that space, and bb3, where the paths meet, reads it and returns. The output lists bb2 first, but the destinations of the br instructions determine execution order. We don't execute every block from top to bottom. LLVM Language Reference

Also, %_0 here is a value pointing to memory, while %0 is the integer read from it. They look almost identical, but that underscore makes them different names. Including the _0 we saw in MIR, it seems safer to follow which instruction creates a value and where it's used, rather than assigning a role based only on its name.

MIR elements don't always correspond one-to-one with LLVM IR elements, either. Function calls may need to handle unwinding after a panic, depending on the settings and the called function. Lowering calls or assert operations can produce multiple LLVM basic blocks from a single MIR basic block.

We saw memory writes and reads in this example, but those aren't guaranteed to remain. Before lowering, rustc also analyzes whether values can be represented directly in SSA form. SSA gives each value name a single definition. Where that form is suitable, rustc can avoid some of the work of building a memory-based representation first and asking LLVM to simplify it afterward.

Lowering from MIR to LLVM IR therefore involves preserving the meaning of Rust operations while building a form LLVM can work with. It takes more than mechanically replacing the notation. Lowering MIR to a Codegen IR — Rust Compiler Development Guide

Generating LLVM IR locally

Save the twice example as twice.rs and the choose example as choose.rs. These commands produce the output used here:

rustc twice.rs --edition 2024 \
  -C opt-level=0 --emit=mir,llvm-ir

rustc choose.rs --crate-type lib --edition 2024 \
  -C opt-level=0 --emit=mir,llvm-ir

Each command produces a .mir file and a .ll file. choose has no main, so it uses --crate-type lib.

The environment was rustc 1.96.0 (ac68faa20 2026-05-25), targeting aarch64-apple-darwin, with LLVM 22.1.2. With optimization kept low under these conditions, we could follow the branches and memory operations. Changing the optimization settings or target environment may change the output's form. I haven't compared compilation time or runtime performance here.

LLVM produces object files, and the process continues

Once the LLVM IR for each CGU is ready, LLVM proceeds with optimization and machine-code generation, producing object files. The two columns in Figure 3 show how this work can proceed in separate units.

After that, the necessary objects and other inputs are linked to produce the output artifact. The kind of LTO in use also affects when optimization happens, but I'll leave the details of linking for another topic. Code generation — Rust Compiler Development Guide

Returning to the original question, “Which code is passed to LLVM, and in what units?”, the connection is becoming clearer: rustc collects the necessary function instances and other items, places them in CGUs, and lowers MIR to LLVM IR for each unit.

Following one if led from handling values to dividing work and deciding how much of a previous compilation could be reused. That division can even affect how fast the generated program runs. There's more going on inside a compiler than carrying out computations. Every topic seems to lead to another thing I want to investigate.

Hmm. Have I finally made it as far as the linker's doorstep...?

For the basics of reading MIR blocks and locals, see Reading Rust's MIR: Following Control Flow and Values.

Reading Rust's MIR: Following Control Flow and Values

Masahiro Morino — Tue, 29 Sep 2026 12:43:23 +0000

I started looking into linkers, but a linker's job isn't specific to Rust. Before getting there, I thought it would be more interesting to see how rustc takes our Rust code toward code generation.

One of the representations we meet along the way is MIR. Apparently, it's used for both borrow checking and code generation. But what exactly is an “intermediate representation,” anyway?

In this article, I'll use a few small examples to explore how to read Rust operations in the compiler's MIR. We'll start with an if expression to get a feel for control flow, look closely at a single assignment to see how values and places are represented, and then move on to a call to Vec::push.

This article is based on the Rust Compiler Development Guide as consulted on September 29, 2026. The examples were checked with Rust 1.96.0. MIR's structure and textual output can change with the compiler version and optimization settings.

What is MIR for?

MIR stands for Mid-level Intermediate Representation. An intermediate representation is a way for a compiler to represent code while processing a program.

The Rust we write has constructs such as if, match, and nested expressions that make programs convenient to write and read. The compiler needs to answer questions like “Where is this value created, and where is it used?” or “After taking this branch, what happens next?” MIR organizes the program into a form that makes types, operations on values, and control flow explicit. That gives the compiler something it can analyze. The MIR — Rust Compiler Development Guide

For example, the expression if flag { x } else { 0 }, which we'll look at shortly, becomes a condition check, operations that set the return value, and a block where the paths meet. Complex expressions are broken down into smaller computations and temporary values. A more uniform representation lets the compiler follow values and control flow without having to handle every variation of the source syntax separately. MIR construction — Rust Compiler Development Guide

One major use of MIR is borrow checking. Remember borrowing? The compiler examines how values are used and borrowed on MIR, and also uses MIR for subsequent optimization and code generation. That doesn't mean one unchanged MIR body survives all the way through compilation. It passes through analysis and transformation stages on its way to the form used for code generation. Overview of the compiler

Let's pause and put these representations side by side.

Representation	Its role in the path we're looking at
Rust source code	Where we write the program using expressions and other language constructs
MIR	Makes Rust types, operations on values, and branches explicit for borrow checking, optimization, and code generation
LLVM IR	Carries the program from rustc into LLVM for further optimization and machine-code generation

There are other compiler representations before MIR, but I'm leaving those out here. I'd like to explore that part of the compiler sometime, too. For now, let's focus on MIR as a representation for analyzing Rust programs and look inside it. Overview of the compiler

In MIR, control flow becomes blocks

One characteristic of MIR is that control flow takes the form of a control-flow graph, or CFG. Operations are grouped into basic blocks, connected by branches and jumps. Each block contains statements, such as assignments, and ends with a terminator, such as a branch, a function call, or a return. The MIR — Rust Compiler Development Guide

Here's a small function:

pub fn choose(flag: bool, x: u32) -> u32 {
    if flag { x } else { 0 }
}

It returns x when flag is true, and 0 when it's false. Compiling this function to MIR with Rust 1.96.0 gives us four basic blocks:

Figure 1: The block structure of the generated MIR. The diagram uses flag, x, and return value in place of the internal local names.

In bb1, setting the return value to x is a statement, and jumping to bb3 is the terminator. bb0 has no statements at all, just the conditional-branch terminator. So a block doesn't have to do some computation before it branches.

A basic block is a unit that executes its statements in order and ends with a terminator. A terminator doesn't necessarily choose between multiple destinations. In this example, goto has one successor, switchInt selects a successor based on a value, and return has no successor block within this function. Thinking of a terminator as the block's exit helps explain why these operations belong to the same category.

Matching the MIR output to the diagram

In the output, arguments and the return value are represented by names such as _1, _2, and _0. Here, _1 is flag, _2 is x, and _0 holds the return value. If we omit the declarations and keep just the blocks, the correspondence with the diagram becomes clear:

bb0: {
    switchInt(copy _1) -> [0: bb2, otherwise: bb1];
}

bb1: {
    _0 = copy _2;
    goto -> bb3;
}

bb2: {
    _0 = const 0_u32;
    goto -> bb3;
}

bb3: {
    return;
}

What was one if expression in the source has become a conditional branch, assignments, and a block where the paths meet. Making control flow explicit gives the compiler a way to analyze each path.

Reading one line of MIR: Places and values

Now that we know how the blocks connect, let's look more closely at the assignment in bb1:

_0 = copy _2;

It simply puts x into the return value, but the pieces have different roles in MIR. The key terms are Local, Place, Operand, and Rvalue. We're starting to collect quite a few terms now.

A Local is storage; a Place identifies a location

A Local is a storage location for something like a function argument, a local variable, or a temporary value. MIR identifies locals by indices, written as _1, for example. These are the locals in choose:

Local	Role in this function	Type
`_0`	The special location for the return value	`u32`
`_1`	The argument `flag`	`bool`
`_2`	The argument `x`	`u32`

Storage here is a compiler-level concept. It doesn't mean that every Local necessarily ends up with its own stack slot in the final machine code. And the same-looking local names can have different roles in different functions, as we'll see below.

A Place is an expression identifying a location to read from or write to. It can refer to an entire Local, a field within it, or a location reached through a reference. For example, if _1 were a tuple, _1 would identify the whole tuple and _1.0 its first field. Yes, this gets a little confusing.

A Local is a unit of storage; a Place specifies a particular location. In _0 = copy _2;, both the destination _0 and the source _2 are Places referring to entire Locals. Key MIR vocabulary

An Operand supplies a value; an Rvalue produces the value to assign

Identifying a location doesn't yet say how we'll use the value stored there. That's where an Operand comes in. Three forms we'll encounter are:

Operand	Meaning
`copy _2`	Copy the value from the specified Place
`move _3`	Move the value from the specified Place
`const 0_u32`	Use a constant value

An Rvalue, meanwhile, is an expression that produces the value to assign. It might add values, create a reference to a Place, or simply use an Operand's value. The R comes from the right-hand side of an assignment.

Here's how the pieces fit together. This is an explanatory tree, not literal rustc output:

Statement: assignment  _0 = copy _2;
├─ Destination Place: _0
└─ Rvalue: use an Operand's value directly
   └─ Operand: copy _2
      └─ Source Place: _2

The printed copy _2 can be a little puzzling: we're looking at an Rvalue when we consider the whole right-hand side, and an Operand when we focus on its input. The Rvalue in this example wraps one Operand and produces that Operand's value unchanged.

The assignment _0 = const 0_u32; in bb2 has the same shape. The input changes from a value stored in a Local to a constant, but we're still assigning an Operand's value to the return-value Place. Assignments, Rvalues, and Operands in MIR

Why is Vec::push a terminator?

Let's use those terms to read another example, also found in the official guide:

fn main() {
    let mut vec = Vec::new();
    vec.push(1);
    vec.push(2);
}

Compiling it to MIR with Rust 1.96.0 produces this block for the first push. I explicitly used -C panic=unwind to inspect the unwinding path:

bb1: {
    _3 = &mut _1;
    _2 = Vec::<i32>::push(move _3, const 1_i32) -> [return: bb2, unwind: bb5];
}

Here, _1 is the Local for vec. _3 is a temporary Local of type &mut Vec<i32>, and _2 receives the () returned by push. Local indices belong to their individual functions, so these _1 and _2 are different from the ones in choose. Watch out for that!

The first line, _3 = &mut _1;, is an assignment statement. On the left, _3 is the destination Place. On the right, &mut _1 is an Rvalue that creates a mutable reference to the Place _1. Our earlier Rvalue simply used an Operand; this one creates a reference value from a Place.

The next line passes the Operands move _3 and const 1_i32 to push. What's being moved is the mutable reference stored in _3, not the Vec itself.

And that whole call is a terminator. Despite its _2 = ... appearance, it isn't an assignment statement. Look at the destinations at the end:

Destination	What happens in this example
`return: bb2`	If the call returns normally, continue to the block containing the next `push`
`unwind: bb5`	If the call unwinds, enter the cleanup block that drops `vec`

The return: bb2 label does not mean that main finishes. It tells us where execution resumes in this function after the called function returns.

In Rust source, vec.push(1); is one line. In MIR, it is split into preparing the reference argument and a call that explicitly includes the paths execution can take afterward. Treating a function call as a terminator lets us follow both normal execution and unwinding cleanup through the connections between blocks. The guide's Vec example

Generating MIR locally

For the first example, I saved the function in choose.rs and ran:

rustc choose.rs --crate-type lib --edition 2024 \
  -C opt-level=0 --emit=mir

This writes the MIR to choose.mir. The environment was rustc 1.96.0 (ac68faa20 2026-05-25), targeting aarch64-apple-darwin, with LLVM 22.1.2. The output shown here comes from those conditions. Different optimization settings, for example, won't necessarily produce the same block structure.

For the Vec example, I saved the code in vec_push.rs and ran:

rustc vec_push.rs --edition 2024 \
  -C opt-level=0 -C panic=unwind --emit=mir

Even with -C opt-level=0, the emitted MIR isn't necessarily the state before all transformations. The StorageLive statements shown in the official guide don't appear in this Vec output. When comparing dumps, we need to pay attention to both the Rust version and the stage of MIR we're looking at. Setting the MIR optimization level with -Z mir-opt-level=0 requires nightly; this article uses output from stable. MIR output and optimization

Reading control flow and values separately

In choose, a single if expression became a block that checks the condition, two blocks that set the return value, and a block that returns. Separating storage locations such as _0 from units of control flow such as bb0 makes the output easier to follow.

Inside each block, Local and Place tell us which location we're dealing with; Operand and Rvalue tell us how values are used and produced. In the Vec::push example, we also saw that a call is itself a terminator, with normal-return and unwinding paths.

MIR represents Rust operations in a form the compiler can analyze. Follow control flow through the connections between blocks, then distinguish places from values within each line. With those two ideas in mind, it becomes easier to find the original Rust operations in what initially looks like a wall of symbols.

So, when we move from MIR toward machine code, what happens to generic type arguments? And how is the code divided into units that can be compiled in parallel? In What Code Does Rust Pass to LLVM? Generics and Codegen Units, I follow the collection of the required code and its partitioning into Codegen Units.

There's a lot to wrap my head around. But that's what makes it interesting.