CUDA en Rust contra CUDA C++: Rust gana por 3,2 % hasta que compruebas que no calculan lo mismo

NVIDIA publicó este mes cuda-oxide, un backend para escribir kernels CUDA en Rust puro. La noticia me llegó por OpenNet, en ruso, y por TabNews, en portugués. En Hacker News el hilo tenía 952 puntos el 16 de septiembre. Leí muchas opiniones sobre el anuncio. Yo quería verlo compilado y medido en mi máquina. Tengo una RTX 4070 Ti SUPER en el escritorio, así que lo instalé, corrí sus ejemplos y escribí el mismo kernel dos veces: una en Rust y otra en CUDA C++.

En Mandelbrot, Rust salió 3,2 % más rápido. Después comprobé que los dos programas no calculan la misma imagen: difieren en 27.510 píxeles, por cómo cada cadena de compilación fusiona multiplicaciones y sumas. Con la aritmética igualada bit a bit, Rust queda 7 % más lento. Los dos números van juntos, y la causa del primero no la pude establecer.

Qué es cuda-oxide

Es un backend de codegen de rustc. Toma las funciones marcadas con #[kernel] y las compila a PTX, el ensamblador intermedio de las GPU de NVIDIA. El host y el dispositivo viven en el mismo archivo Rust, y todo se maneja con cargo oxide build y cargo oxide run. La licencia es Apache 2.0. El README no se anda con rodeos sobre su estado: es alpha, y dice textualmente “you should expect bugs, incomplete features, and API breakage”.

El anuncio de NVIDIA presenta dos caminos. El otro se llama cutile-rs. No lo probé, y en este post no digo nada sobre él.

Este es el kernel que medí, en 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;
    }
}

Y el mismo en 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;
}

Las dos fuentes son equivalentes operación por operación en punto flotante. En los enteros no: son u32 en Rust e int en C++, y no probé igualarlos. Guarda los dos datos, porque más abajo importan.

Instalación, con los tropiezos

git clone https://github.com/NVlabs/cuda-oxide
cd cuda-oxide                             # rustup baja solo el nightly que fija el repo
cargo build --release                     # falló a los 37 s, ver abajo
sudo dnf install libcurand-devel-13-1     # 256 MB entre biblioteca y cabeceras
cargo build --release                     # ahora sí: 25 s
cargo oxide run host_closure              # el ejemplo del README

• No sirve con Rust estable. Exige nightly-2026-08-28 con rustc-dev, rust-src y llvm-tools. Lo fija el rust-toolchain.toml del repositorio, y rustup lo descarga solo al entrar al directorio.
• El primer cargo build --release falló a los 37 s con fatal error: 'curand.h' file not found. El toolkit de CUDA de mi máquina traía nvcc, pero no las cabeceras de cuRAND. Lo arreglé con dnf install libcurand-devel-13-1, que son 256 MB entre biblioteca y cabeceras. También necesita clang.
• Con eso el backend compila en 25 s. cargo oxide run host_closure, el ejemplo del README, compila en 1 minuto y pasa sus 13 pruebas en la GPU real. En el commit que medí hay 230 directorios de ejemplos; el README dice “190+”.
• cargo oxide run detecta la arquitectura de la tarjeta, en mi caso sm_89. cargo oxide build sin --arch apunta a sm_80 y lo avisa. En los dos casos el binario lleva el PTX como texto y el driver lo compila con JIT al cargar el módulo. Lo único que lo evita es --materialize-cubin.
• Lo que no logré: compilar Rust por adelantado (AOT). --materialize-cubin necesita libnvJitLink.so, que falta en mi instalación, y no la agregué para esta medición. Todo lo de Rust en este post pasó por el JIT del driver.

Cómo medí

La máquina: Fedora 43 con kernel 7.2.4, Ryzen 7 7800X3D y 30 GB de RAM. La GPU es una RTX 4070 Ti SUPER (Ada, sm_89, 16 GB, unos 672 GB/s teóricos de ancho de banda), con driver 595.91.07, límite de potencia de 285 W y persistence mode apagado. CUDA 13.1 (nvcc V13.1.115) y gcc 15.3.1. De cuda-oxide usé el commit b9847e9, con rustc 1.100.0-nightly, cargo-oxide 0.2.1 y cuda-core 0.3.1.

Dos kernels. SAXPY con 2^24 elementos, que está limitado por memoria. Y Mandelbrot de 4096×4096 con 500 iteraciones, limitado por cómputo.

El método es igual en los dos lenguajes. Bloques de 256 hilos, fijados a mano e impresos en el log de cada corrida. Calentamiento de al menos 300 ms, después 300 repeticiones cronometradas con eventos CUDA. Diez rondas intercaladas, rotando el orden en cada una, lo que deja n = 10 por variante. Y un control A/A: el mismo binario de C++ medido contra una copia de sí mismo, para saber cuánto ruido tiene el banco.

Estas son las líneas de compilación de Mandelbrot, con y sin contracción FMA. El ejemplo de Rust vive dentro del clon, en 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: el control, y un empate

0,324 ms por iteración en ambos lenguajes. Son unos 621 GB/s, el 92 % del ancho de banda teórico de la tarjeta. Ahí decide el ancho de banda de la memoria de la tarjeta, y al compilador le queda poco que aportar.

Rust salió entre +0,01 y +0,02 % más lento, que son unos 30 a 55 ns por lanzamiento. El signo se repite en las dos variantes de Rust, pero la magnitud es irrelevante: el control A/A varía ±0,03 %. Lo cuento como empate, sin fingir que la diferencia es cero exacto.

Mandelbrot con las opciones por defecto: Rust 3,2 % más rápido

Con las opciones por defecto, Rust tardó 0,884 ms por lanzamiento y C++ 0,915 ms. La diferencia pareada, ronda contra ronda, es −3,25 % (rango −3,37 a −2,99), y Rust quedó adelante en 10 de 10 rondas. La mediana del control A/A aquí es −0,02 %, y el signo no cambió en ninguna ronda. Ruido no es.

Desconfié del instrumento y no era eso. Da lo mismo con eventos que con reloj de pared: sobre la misma ventana, el reloj de pared sesga +0,003 % en Mandelbrot y +0,007 % en SAXPY. El JIT cuesta unos 13 ms y queda fuera de la ventana cronometrada. Tampoco cambia al compilar para sm_80 con JIT o para sm_89, y en C++ no hay diferencia medible entre AOT y JIT.

Lo que sí mueve la brecha es el tamaño del bloque: −3,4 % con 128 hilos, −3,2 % con 256 y −2,3 % con 512, en un control corto de n = 3. Fijar el bloque en los dos lados importaba.

Hasta aquí tenía un titular. Me faltaba comprobar algo que en la primera pasada di por hecho.

Los dos programas no calculan lo mismo

Cada binario imprime la suma de los 16.777.216 valores de salida y un hash del búfer completo. No coinciden.

Para entender por qué hay que saber qué es una FMA. Es una instrucción de fused multiply-add: calcula a*b + c de una sola vez y redondea solo al final. Hecho en dos pasos se redondea dos veces, primero el producto y después la suma, y el último bit puede salir distinto.

Un último bit parece nada. En Mandelbrot no lo es, porque cada iteración alimenta a la siguiente y cerca del borde del conjunto una diferencia mínima decide si el punto escapa ahora o mucho después. El resultado, con las opciones por defecto:

• 27.510 píxeles distintos de 16.777.216, el 0,164 %.
• Hasta 392 iteraciones de diferencia en un mismo píxel.
• 3.484 píxeles cambian de clase: en una imagen están dentro del conjunto y en la otra fuera.

En punto flotante las dos fuentes piden lo mismo. Lo que cambia es qué multiplicaciones y sumas fusiona cada cadena de compilación.

Píldora violeta: en esa etapa se contrae una FMA del bucle interno.

La tercera fusión de Rust (zx*zx − zy*zy) no está en el PTX que escribe cuda-oxide: la agrega la etapa de NVIDIA que convierte el PTX en código de máquina. Por eso contar FMA en el archivo intermedio me dio la imagen equivocada.

Ninguna de las dos imágenes es la de float32 estricto. La de nvcc se aparta en 27.592 píxeles y la de Rust en 31.606.

Para no fiarme solo de la GPU escribí una emulación en CPU del mismo bucle, con los tres patrones de fusión. Reproduce los tres hashes exactos.

La emulación también muestra qué fusión pesa. La de zy la hacen las dos cadenas. La de la prueba de escape, zx*zx + zy*zy, no mueve ningún píxel, así que los hashes no sirven para validarla. Toda la diferencia entre la imagen de Rust y la de nvcc viene de una sola fusión: zx*zx − zy*zy.

Hay una explicación tentadora que descarté con números: que Rust gane porque hace menos trabajo. Sobre esos 27.510 píxeles distintos, el neto es de 8.912 iteraciones menos para Rust, con diferencias que se cancelan entre píxeles, sobre un total de 1,49 × 10⁹. Es un 0,0006 % menos de iteraciones. Eso no explica un 3,2 %.

cuda-oxide contrae a propósito. Imita a nvcc, y tiene documentado el flag --no-fmad para apagarlo. Pero es una desviación semántica que conviene conocer: rustc en CPU nunca contrae a*b + c por su cuenta, y quien quiere una FMA la pide con f32::mul_add. En este bucle, además, termina contrayendo más que nvcc.

Con la aritmética igualada, Rust pierde por 7 %

Apagué la contracción en los dos lados: -fmad=false en nvcc y --no-fmad en cuda-oxide. Las salidas pasan a ser idénticas bit a bit. Y el orden se invierte: Rust queda +7,35 % más lento en la medición pareada (rango +6,80 a +7,38), con 0 de 10 rondas a su favor.

Apagar FMA le cuesta +4,10 % a C++ y +15,55 % a Rust.

← Rust más rápidoRust más lento →

Opciones por defectolas dos cadenas fusionan FMA, cada una a su manera

Rust −3,25 %rango −3,37 a −2,99 %Rust adelante en 10 de 10 rondas27.510 píxeles distintos

FMA apagada en ambos lados-fmad=false en nvcc, --no-fmad en cuda-oxide

Rust +7,35 %rango +6,80 a +7,38 %Rust adelante en 0 de 10 rondassalidas idénticas bit a bit

−5 %empate+5 %+10 %

Ninguno de los dos números va solo. El primero compara dos programas que no calculan la misma imagen; el segundo compara resultados idénticos, y ahí Rust pierde. No hay una configuración medida donde Rust gane y las salidas coincidan.

No existe una configuración medida donde Rust gane y los resultados coincidan. Si alguien cita el −3,25 % de este post como “Rust contra C++” a secas, lo está citando mal: compara dos programas que no producen la misma imagen.

La causa: no la establecí

Fui a mirar el código que de verdad ejecuta la GPU. Para C++ desensamblé la imagen AOT de nvcc. Para Rust capturé la imagen que el JIT del driver produjo para el mismo binario que medí, y desensamblé esa.

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 instrucciones · 4 sumas + 2 multiplicaciones + 1 FMA · un salto combinado

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 instrucciones · 2 sumas + 1 multiplicación + 3 FMA · dos saltos condicionales

FMA: multiplica y suma con un solo redondeo

El mismo largo, distinta aritmética. Con la contracción apagada los bucles pasan a 12 instrucciones en C++ y 13 en Rust. El desensamblador que usé es de 2024, anterior al toolkit; su decodificación cuadra con las sumas de control, pero es una salvedad.

Los dos bucles internos tienen 11 instrucciones por iteración. C++ hace 4 sumas, 2 multiplicaciones y 1 FMA, que son 7 instrucciones de punto flotante, y cierra con un único salto combinado. Rust hace 2 sumas, 1 multiplicación y 3 FMA, 6 instrucciones, con dos saltos condicionales. Contando cada FMA como dos operaciones aritméticas, una multiplicación y una suma, son 8 en C++ y 9 en Rust, así que un conteo de instrucciones no permite leer el 3 % en el bucle. Dos de esas fusiones ya vienen en el PTX que emite cuda-oxide. La tercera, zx*zx − zy*zy, no está en ese PTX. La agrega la etapa de NVIDIA que lo convierte en código de máquina: el JIT del driver, y también ptxas, el ensamblador offline del toolkit, que produce el mismo bucle.

¿Gana Rust por tener una instrucción flotante menos? Es mi hipótesis y no la comprobé. El experimento de apagar FMA no sirve para aislarla, porque cambia dos cosas a la vez: la aritmética y la forma del bucle, que pasa a 12 instrucciones en C++ y 13 en Rust. Sé que el orden se invierte al quitar la contracción. No sé por qué.

Hay otra variable que no controlé y que ya venía en la fuente: los enteros. Son u32 en Rust e int en C++. En el desensamblado, C++ cierra el bucle con n >= iter_max y Rust con n != iter_max. No probé la versión de C++ con unsigned.

Lo que tuve que retractar

En la primera pasada conté instrucciones fma en el PTX intermedio, 2 en Rust contra 1 en C++, y lo di por explicación. Al revisarlo vi que el total de operaciones era el mismo y que el PTX no es lo que ejecuta la GPU: la etapa siguiente todavía lo transforma, como muestra la tercera fusión.

El otro error fue peor. El chequeo de equivalencia de esa pasada era vacío. Comparaba dos píxeles que escapan sin ejercer redondeo alguno, así que daba “iguales” con cualquier aritmética. Un chequeo que no puede fallar no comprueba nada, y con él habría publicado el 3,2 % como una victoria limpia.

Condiciones que no controlé

Las condiciones no fueron idénticas entre variantes, y prefiero decirlo.

• Tras las corridas de C++ el reloj gráfico marcó un escalón menos, 2820 en vez de 2835 MHz, en 12 de 40 casos. Tras las de Rust, nunca. C++ consumió unos 6 W más. No medí el reloj dentro de la ventana cronometrada.
• La temperatura subió de 39 a 73 °C durante la sesión. Razonado, no verificado: un escalón de reloj da a lo sumo 0,5 %, y la brecha se mantuvo en −3,0 % con la GPU fría.
• ComfyUI estuvo residente en la GPU con 210 MiB, aparentemente inactivo. Lo infiero, no lo medí.
• Mi instalación de CUDA 13.1 no traía cuobjdump ni nvdisasm. El desensamblador que usé es de 2024, anterior al toolkit. Su decodificación cuadra con los hashes, pero queda como salvedad.

Mi opinión

Me sorprendió lo poco que costó. Un tropiezo con cabeceras, 25 s de compilación, y tenía kernels de Rust corriendo en la GPU con cargo. El host y el dispositivo en un mismo archivo, con los tipos de Rust de punta a punta, se siente mejor que mantener un .cu aparte con su FFI. Y en el kernel limitado por memoria no hubo nada que discutir: empate.

Lo que no: es alpha y se nota en los bordes. Nightly fijado a una fecha, AOT que no pude armar, y una decisión de aritmética que el lenguaje no toma en ningún otro lado. Quien viene de Rust en CPU espera que a*b + c redondee dos veces. Aquí no, salvo que pases --no-fmad, y pasarlo costó +15,55 % en mi kernel.

Sobre el rendimiento, lo único que sostengo es esto: en un kernel de cómputo, cuda-oxide quedó entre un 3,2 % por debajo y un 7 % por arriba del tiempo de nvcc, según qué aritmética aceptes. Para una primera versión pública me parece un buen punto de partida. Decir que Rust es más rápido que C++ en GPU sería inventar.

Cuándo lo usaría

• Sí: prototipos y herramientas propias donde el host ya está en Rust; para aprender CUDA sin salir de cargo; kernels como mi SAXPY, donde manda el ancho de banda de la memoria.
• Sí, con cuidado: cómputo numérico donde importa el último bit. Decide desde el principio si vas con FMA o con --no-fmad, y compara contra una referencia con píxeles que de verdad ejerzan el redondeo.
• Todavía no: producción. Tampoco si necesitas reproducir bit a bit una salida de CUDA C++ con las opciones por defecto: las dos cadenas contraen distinto y no vi un ajuste que las iguale sin apagar FMA en ambas.

Fuentes

• Introducing CUDA Rust: Two Tracks for Writing GPU Kernels, blog de desarrolladores de NVIDIA (inglés). Fuente original del anuncio.
• NVlabs/cuda-oxide, el repositorio, con el README que lo declara alpha. Medí el commit b9847e9.
• README de cargo-oxide, donde está documentado --no-fmad.
• NVlabs/cutile-rs, el otro camino del anuncio, que no probé.
• Floating Point and IEEE 754 Compliance for NVIDIA GPUs, el documento de NVIDIA que explica por qué una FMA redondea una sola vez.
• Documentación de nvcc, por la opción -fmad.
• f32::mul_add en la biblioteca estándar de Rust.
• El hilo en Hacker News.