CUDA in Rust vs CUDA C++: Rust wins by 3.2% until you check that they do not compute the same thing

This month NVIDIA released cuda-oxide, a backend for writing CUDA kernels in pure Rust. The news reached me through OpenNet, in Russian, and TabNews, in Portuguese. On Hacker News the thread had 952 points on September 16. I read plenty of opinions about the announcement. I wanted to see it built and measured on my own machine. There is an RTX 4070 Ti SUPER on my desk, so I installed it, ran its examples and wrote the same kernel twice: once in Rust and once in CUDA C++.

On Mandelbrot, Rust came out 3.2% faster. Then I checked, and the two programs do not compute the same image: they differ in 27,510 pixels, because of how each toolchain fuses multiplies and adds. With the arithmetic matched bit for bit, Rust ends up 7% slower. The two numbers belong together, and I could not establish the cause of the first one.

What cuda-oxide is

It is a codegen backend for rustc. It takes the functions marked #[kernel] and compiles them to PTX, the intermediate assembly of NVIDIA GPUs. Host and device live in the same Rust file, and everything is driven through cargo oxide build and cargo oxide run. The license is Apache 2.0. The README is blunt about its state: it is alpha, and it says “you should expect bugs, incomplete features, and API breakage”.

NVIDIA’s announcement presents two tracks. The other one is called cutile-rs. I did not try it, and this post says nothing about it.

This is the kernel I measured, in Rust:

#[kernel]
pub fn mandel(w: u32, h: u32, iter_max: u32, mut salida: DisjointSlice<i32>) {
    let idx = thread::index_1d();
    let i = idx.get() as u32;
    if let Some(celda) = salida.get_mut(idx) {
        let cx = -2.0f32 + 3.0f32 * ((i % w) as f32) / (w as f32);
        let cy = -1.5f32 + 3.0f32 * ((i / w) as f32) / (h as f32);
        let (mut zx, mut zy) = (0.0f32, 0.0f32);
        let mut n: u32 = 0;
        while n < iter_max && zx * zx + zy * zy <= 4.0f32 {
            let t = zx * zx - zy * zy + cx;
            zy = 2.0f32 * zx * zy + cy;
            zx = t;
            n += 1;
        }
        *celda = n as i32;
    }
}

And the same one in CUDA C++:

__global__ void mandel(int w, int h, int iter_max, int* salida) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= w * h) return;
    float cx = -2.0f + 3.0f * (float)(i % w) / (float)w;
    float cy = -1.5f + 3.0f * (float)(i / w) / (float)h;
    float zx = 0.0f, zy = 0.0f;
    int n = 0;
    while (n < iter_max && zx * zx + zy * zy <= 4.0f) {
        float t = zx * zx - zy * zy + cx;
        zy = 2.0f * zx * zy + cy;
        zx = t;
        n++;
    }
    salida[i] = n;
}

The two sources are equivalent operation by operation in floating point. The integers are not: they are u32 in Rust and int in C++, and I did not try matching them. Keep both facts in mind, because they matter further down.

Installing it, stumbles included

git clone https://github.com/NVlabs/cuda-oxide
cd cuda-oxide                             # rustup fetches the nightly the repo pins
cargo build --release                     # failed after 37 s, see below
sudo dnf install libcurand-devel-13-1     # 256 MB of library plus headers
cargo build --release                     # now it works: 25 s
cargo oxide run host_closure              # the example from the README

• It does not work on stable Rust. It requires nightly-2026-08-28 with rustc-dev, rust-src and llvm-tools. The repository’s rust-toolchain.toml pins it, and rustup downloads it on its own when you enter the directory.
• The first cargo build --release failed after 37 s with fatal error: 'curand.h' file not found. The CUDA toolkit on my machine shipped nvcc but not the cuRAND headers. I fixed it with dnf install libcurand-devel-13-1, which is 256 MB of library plus headers. It also needs clang.
• With that the backend builds in 25 s. cargo oxide run host_closure, the README example, builds in 1 minute and passes its 13 tests on the real GPU. At the commit I measured there are 230 example directories; the README says “190+”.
• cargo oxide run detects the card’s architecture, sm_89 in my case. cargo oxide build without --arch targets sm_80 and says so. Either way the binary carries the PTX as text and the driver JIT-compiles it when the module loads. The only thing that avoids it is --materialize-cubin.
• What I did not manage: ahead-of-time (AOT) compilation for Rust. --materialize-cubin needs libnvJitLink.so, which is missing from my install, and I did not add it for this measurement. Everything Rust in this post went through the driver’s JIT.

How I measured

The machine: Fedora 43 on kernel 7.2.4, a Ryzen 7 7800X3D and 30 GB of RAM. The GPU is an RTX 4070 Ti SUPER (Ada, sm_89, 16 GB, about 672 GB/s of theoretical bandwidth), driver 595.91.07, a 285 W power limit and persistence mode off. CUDA 13.1 (nvcc V13.1.115) and gcc 15.3.1. For cuda-oxide I used commit b9847e9, with rustc 1.100.0-nightly, cargo-oxide 0.2.1 and cuda-core 0.3.1.

Two kernels. SAXPY over 2^24 elements, which is memory-bound. And a 4096×4096 Mandelbrot with 500 iterations, which is compute-bound.

The method is the same in both languages. Blocks of 256 threads, set by hand and printed in the log of every run. A warm-up of at least 300 ms, then 300 repetitions timed with CUDA events. Ten interleaved rounds, rotating the order in each one, which gives n = 10 per variant. And an A/A control: the same C++ binary measured against a copy of itself, to learn how noisy the bench is.

These are the build lines for Mandelbrot, with and without FMA contraction. The Rust example lives inside the clone, under crates/rustc-codegen-cuda/examples/mandel_b2.

nvcc -O3 -arch=sm_89 -o mandel_cpp mandel.cu
nvcc -O3 -arch=sm_89 -fmad=false -o mandel_cpp_nofmad mandel.cu
cargo oxide build mandel_b2 --arch sm_89
cargo oxide build mandel_b2 --arch sm_89 --no-fmad

SAXPY: the control, and a tie

0.324 ms per iteration in both languages. That is about 621 GB/s, 92% of the card’s theoretical bandwidth. The card’s memory bandwidth decides there, and the compiler has little left to add.

Rust came out between +0.01 and +0.02% slower, which is about 30 to 55 ns per launch. The sign repeats across both Rust variants, but the size is irrelevant: the A/A control varies by ±0.03%. I call it a tie, without pretending the difference is exactly zero.

Mandelbrot with default options: Rust 3.2% faster

With default options, Rust took 0.884 ms per launch and C++ 0.915 ms. The paired difference, round against round, is −3.25% (range −3.37 to −2.99), and Rust was ahead in 10 of 10 rounds. The median of the A/A control here is −0.02%, and the sign did not change in any round. It is not noise.

I suspected the instrument and that was not it. Events and wall clock agree: over the same window, the wall clock is biased by +0.003% on Mandelbrot and +0.007% on SAXPY. The JIT costs about 13 ms and stays outside the timed window. Nothing changes between building for sm_80 with JIT and for sm_89 either, and in C++ there is no measurable difference between AOT and JIT.

What does move the gap is the block size: −3.4% with 128 threads, −3.2% with 256 and −2.3% with 512, in a short control with n = 3. Pinning the block size on both sides mattered.

At this point I had a headline. I still had to check something I took for granted in my first pass.

The two programs do not compute the same thing

Each binary prints the sum of the 16,777,216 output values and a hash of the whole buffer. They do not match.

To see why, you need to know what an FMA is. It is a fused multiply-add instruction: it computes a*b + c in one go and rounds only at the end. Done in two steps it rounds twice, first the product and then the sum, and the last bit can come out different.

One last bit sounds like nothing. In Mandelbrot it is not, because each iteration feeds the next one, and near the border of the set a tiny difference decides whether the point escapes now or much later. The result, with default options:

• 27,510 pixels differ out of 16,777,216, or 0.164%.
• Up to 392 iterations of difference in a single pixel.
• 3,484 pixels change class: inside the set in one image and outside in the other.

In floating point the two sources ask for the same thing. What changes is which multiplies and adds each toolchain fuses.

Violet pill: an FMA contraction of the inner loop happens at this stage.

The third Rust fusion (zx*zx − zy*zy) is not in the PTX that cuda-oxide writes: the NVIDIA stage that turns the PTX into machine code adds it. That is why counting FMAs in the intermediate file gave me the wrong picture.

Neither image is the strict float32 one. The nvcc image departs from it in 27,592 pixels and the Rust image in 31,606.

So as not to trust the GPU alone, I wrote a CPU emulation of the same loop with the three fusion patterns. It reproduces all three hashes exactly.

The emulation also shows which fusion matters. The zy one is done by both chains. The one in the escape test, zx*zx + zy*zy, moves no pixel, so the hashes cannot validate it. The whole difference between the Rust image and the nvcc image comes from a single fusion: zx*zx − zy*zy.

There is a tempting explanation that the numbers rule out: that Rust wins because it does less work. Across those 27,510 differing pixels, the net is 8,912 fewer iterations for Rust, with differences that cancel out between pixels, over a total of 1.49 × 10⁹. That is 0.0006% fewer iterations. It does not explain 3.2%.

cuda-oxide contracts on purpose. It mirrors nvcc, and it has the --no-fmad flag documented to turn it off. But it is a semantic deviation worth knowing about: rustc on CPU never contracts a*b + c on its own, and whoever wants an FMA asks for it with f32::mul_add. In this loop, on top of that, it ends up contracting more than nvcc.

With the arithmetic matched, Rust loses by 7%

I turned contraction off on both sides: -fmad=false in nvcc and --no-fmad in cuda-oxide. The outputs become bit-identical. And the order flips: Rust is +7.35% slower in the paired measurement (range +6.80 to +7.38), with 0 of 10 rounds in its favor.

Turning FMA off costs C++ +4.10% and Rust +15.55%.

← Rust fasterRust slower →

Default optionsboth chains fuse FMA, each its own way

Rust −3.25 %range −3.37 to −2.99 %Rust ahead in 10 of 10 rounds27,510 pixels differ

FMA off on both sides-fmad=false in nvcc, --no-fmad in cuda-oxide

Rust +7.35 %range +6.80 to +7.38 %Rust ahead in 0 of 10 roundsbit-identical outputs

−5 %tie+5 %+10 %

Neither number stands alone. The first compares two programs that do not compute the same image; the second compares identical results, and there Rust loses. There is no measured configuration where Rust wins and the outputs match.

There is no measured configuration where Rust wins and the results match. Anyone quoting the −3.25% from this post as plain “Rust against C++” is quoting it wrong: it compares two programs that do not produce the same image.

The cause: I did not establish it

I went to look at the code the GPU actually runs. For C++ I disassembled the AOT image from nvcc. For Rust I captured the image the driver’s JIT produced for the very binary I timed, and disassembled that.

CUDA C++ · nvcc

1. FADDR0 = R3 + (−R0)
2. FADDR3 = R6 + R6
3. IADD3n++
4. FADDR6 = cx + R0
5. FFMAR7 = R3*R7 + cy
6. ISETPn >= iter_max
7. FMULR3 = zx*zx
8. FMULR0 = zy*zy
9. FADDR8 = R3 + R0
10. FSETPR8 <= 4
11. BRA@!P0 P1, loop

11 instructions · 4 adds + 2 multiplies + 1 FMA · one combined branch

Rust · cuda-oxide

1. FMULR7 = zy*zy
2. FFMAR9 = zx*zx + R7
3. FSETPR9 > 4
4. BRA@P0 exit
5. IADD3n++
6. FFMAR7 = zx*zx + (−R7)
7. FADDR9 = zx + zx
8. ISETPn != iter_max
9. FADDzx = cx + R7
10. FFMAzy = R9*zy + cy
11. BRA@P0 loop

11 instructions · 2 adds + 1 multiply + 3 FMA · two conditional branches

FMA: multiplies and adds with a single rounding

Same length, different arithmetic. With contraction off the loops grow to 12 instructions in C++ and 13 in Rust. The disassembler I used is from 2024, older than the toolkit; its decoding matches the checksums, but it is a caveat.

Both inner loops have 11 instructions per iteration. C++ does 4 adds, 2 multiplies and 1 FMA, which is 7 floating-point instructions, and closes with a single combined branch. Rust does 2 adds, 1 multiply and 3 FMAs, 6 instructions, with two conditional branches. Counting each FMA as two arithmetic operations, a multiply and an add, that is 8 for C++ and 9 for Rust, so an instruction count does not let you read the 3% off the loop. Two of those fusions already come in the PTX that cuda-oxide emits. The third one, zx*zx − zy*zy, is not in that PTX. It is added by the NVIDIA stage that turns it into machine code: the driver’s JIT, and also ptxas, the toolkit’s offline assembler, which produces the same loop.

Does Rust win because it has one floating-point instruction fewer? That is my hypothesis and I did not test it. The FMA-off experiment cannot isolate it, because it changes two things at once: the arithmetic and the shape of the loop, which grows to 12 instructions in C++ and 13 in Rust. I know the order flips when contraction goes away. I do not know why.

There is another variable I did not control, and it was already in the source: the integers. They are u32 in Rust and int in C++. In the disassembly, C++ closes the loop with n >= iter_max and Rust with n != iter_max. I did not try the C++ version with unsigned.

What I had to retract

In my first pass I counted fma instructions in the intermediate PTX, 2 in Rust against 1 in C++, and took that as the explanation. On review, the total number of operations was the same, and the PTX is not what the GPU runs: the next stage still transforms it, as the third fusion shows.

The other mistake was worse. The equivalence check in that pass was empty. It compared two pixels that escape without exercising any rounding at all, so it said “equal” under any arithmetic. A check that cannot fail checks nothing, and with it I would have published the 3.2% as a clean win.

Conditions I did not control

Conditions were not identical across variants, and I would rather say so.

• After the C++ runs the graphics clock read one step lower, 2820 instead of 2835 MHz, in 12 of 40 cases. After the Rust runs, never. C++ drew about 6 W more. I did not measure the clock inside the timed window.
• Temperature rose from 39 to 73 °C during the session. Reasoned, not verified: one clock step is worth 0.5% at most, and the gap held at −3.0% with the GPU cold.
• ComfyUI sat resident on the GPU with 210 MiB, apparently idle. I infer that; I did not measure it.
• My CUDA 13.1 installation did not contain cuobjdump or nvdisasm. The disassembler I used is from 2024, older than the toolkit. Its decoding matches the hashes, but it stays as a caveat.

My take

I was surprised by how little it took. One stumble over headers, 25 s of compiling, and I had Rust kernels running on the GPU through cargo. Host and device in one file, with Rust types end to end, feels better than keeping a separate .cu with its FFI. And on the memory-bound kernel there was nothing to argue about: a tie.

The downside: it is alpha and it shows at the edges. A nightly pinned to a date, AOT I could not get working, and an arithmetic decision the language makes nowhere else. Someone coming from Rust on CPU expects a*b + c to round twice. Not here, unless you pass --no-fmad, and passing it cost +15.55% on my kernel.

On performance, the only thing I stand behind is this: on a compute kernel, cuda-oxide landed between 3.2% under and 7% over nvcc’s time, depending on which arithmetic you accept. For a first public release that looks like a good starting point to me. Saying Rust is faster than C++ on the GPU would be making things up.

When I would use it

• Yes: prototypes and in-house tools where the host is already Rust; learning CUDA without leaving cargo; kernels like my SAXPY, where memory bandwidth rules.
• Yes, with care: numerical work where the last bit matters. Decide from the start whether you go with FMA or with --no-fmad, and compare against a reference using pixels that really exercise rounding.
• Not yet: production. Nor if you need to reproduce a default-options CUDA C++ output bit for bit: the two chains contract differently, and I saw no setting that matches them without turning FMA off on both.

Sources

• Introducing CUDA Rust: Two Tracks for Writing GPU Kernels, NVIDIA developer blog. The original announcement.
• NVlabs/cuda-oxide, the repository, with the README that calls it alpha. I measured commit b9847e9.
• cargo-oxide README, where --no-fmad is documented.
• NVlabs/cutile-rs, the other track of the announcement, which I did not try.
• Floating Point and IEEE 754 Compliance for NVIDIA GPUs, NVIDIA’s paper on why an FMA rounds only once.
• nvcc documentation, for the -fmad option.
• f32::mul_add in the Rust standard library.
• The Hacker News thread.