NVIDIA publicó el 31 de octubre de 2026 la versión 0.3.0 de cuda-oxide, el compilador que permite escribir kernels CUDA en Rust, y con ella deja atrás la dependencia de las versiones nightly del lenguaje: el proyecto se compila ahora con Rust estable 1.99.1 y comparte un mismo runtime de host con cutile-rs dentro del monorepositorio NVIDIA/cuda-rust.
Las claves de CUDA Oxide 0.3 en 30 segundos
- NVIDIA publicó CUDA Oxide 0.3 el 31/10/2026, con casi 600 commits de más de 45 colaboradores desde la 0.2.1.
- El compilador SIMT deja las versiones nightly y fija Rust estable 1.99.1, aunque sigue necesitando rustc-dev, herramientas LLVM y CUDA 13.0 o superior.
- cuda-oxide y cutile-rs comparten ya un runtime de host con buffers, streams, eventos y CUDA graphs.
- cute-rs, todavía en vista previa, lleva CuTe a Rust solo para Blackwell sm_100a y sm_120a.
- NVIDIA mantiene el proyecto como experimental.
Según las notas de la versión en GitHub, es la mayor entrega del proyecto hasta ahora. NVIDIA presentó CUDA Rust en septiembre como una plataforma con dos formas de escribir kernels, y en revistacloud.com ya se contó entonces cómo NVIDIA llevaba CUDA a Rust sin salir del lenguaje. Aquella primera etapa exigía una toolchain nightly concreta para cuda-oxide. La 0.3 cambia esa base, aunque no la simplifica del todo.
El compilador sigue fijando su versión de Rust en un fichero rust-toolchain.toml y necesita los componentes rustc-dev, rust-src y llvm-tools. El motivo es técnico: cuda-oxide es un backend de generación de código para rustc y se apoya en componentes internos del compilador (la biblioteca librustc_driver y rustc_public), que no tienen garantías de estabilidad entre versiones. La herramienta cargo-oxide gestiona esa toolchain fijada. La diferencia práctica es que ya no hay que seguir un nightly que cambia cada día.
SIMT y Tile: dos compiladores sobre un runtime común
El repositorio NVIDIA/cuda-rust agrupa dos modelos de programación. cuda-oxide sigue el modelo SIMT (Single Instruction, Multiple Threads, una instrucción ejecutada por muchos hilos), el clásico de CUDA C++: el desarrollador escribe lo que hace un hilo y controla bloques, warps y memoria compartida. Compila por adelantado de Rust a PTX (Parallel Thread Execution, el lenguaje intermedio de las GPU de NVIDIA) pasando por la representación MIR de rustc, el marco de IR Pliron y LLVM.
cutile-rs implementa el modelo CUDA Tile, que NVIDIA ya ofrece en C++ y Python. Aquí el kernel trabaja con bloques de datos completos (tiles) y el compilador decide cómo se reparten en cada arquitectura; el código se compila en el primer lanzamiento hacia Tile IR y binario cubin. Este medio explicó en su día cómo funciona CUDA Tile. cutile-rs no necesita la toolchain fijada: funciona con Rust estable 1.89 o posterior, y su documentación recomienda CUDA 13.3, que es imprescindible para el empaquetado FP4 y las MMA con escalado por bloques.
La novedad de la 0.3 es que ambos comparten el runtime de host formado por cuda-bindings, cuda-core, cuda-core-derive y cuda-async (versión 0.5.0). Las notas de la versión describen un único DeviceBuffer, Stream, Event, CudaGraph y DeviceOp asíncrono para los dos modelos, de modo que un kernel Tile y uno SIMT pueden trabajar en el mismo stream sobre los mismos buffers. También añade gestión de memoria virtual, acceso entre GPU, multicast por NVLink Switch y prioridades de stream, además de un atributo #[tile_fn] para exportar funciones SIMT como bibliotecas que un kernel Tile puede llamar.
cute-rs: CUTLASS en Rust, solo para dos variantes de Blackwell
La pieza más llamativa es cute-rs-preview, una implementación experimental de CuTe, la capa de layouts de CUTLASS (la biblioteca de plantillas de NVIDIA para álgebra lineal en GPU). Compila por adelantado mediante el compilador fijado de CUTLASS 4.7 y cubre, según su propia matriz de compatibilidad, un subconjunto validado de CuTeDSL 4.7, sin paridad de funciones.
Los objetivos son estrictos. En la 0.3.0 solo admite sm_100a (Blackwell de centro de datos, clase B200) y sm_120a (GeForce RTX 50). Ampere, Ada y Hopper (sm_90a) se rechazan antes de compilar, igual que otras variantes de Blackwell. La matriz indica además que no hay soporte de wgmma, la instrucción de Tensor Cores propia de Hopper, ni de kernels agrupados, MoE, entradas mixtas o FP8. En la rama principal, después del lanzamiento, ya figura también sm_103a (clase B300), pero no forma parte de la versión 0.3.0.
Lo que sí ofrece es expresar en el sistema de tipos de Rust los layouts, los descriptores TMA (Tensor Memory Accelerator, el motor de copias asíncronas de memoria), los pipelines productor-consumidor y las MMA de Tensor Cores, y comprobarlos antes de generar código. La versión incluye 14 kernels de ejemplo y dos modelos completos que funcionan en una RTX 5090: Qwen3-8B de 36 capas en BF16 (bfloat16, coma flotante de 16 bits) y NVFP4 (el formato de 4 bits con escalado por bloques de NVIDIA), y Qwen3-TTS 1.7B. En NVFP4 los pesos ocupan 3,64 GiB frente a 12,94 GiB en BF16, 3,56 veces menos.
Depuración con cuda-gdb, pruebas y requisitos de instalación
Para quien tenga que mantener kernels, la 0.3 trae depuración DWARF (el formato estándar de información de depuración) completa con cuda-gdb mediante las opciones --device-debug y --lineinfo, que conservan ubicaciones de código, closures, variantes de enum y memoria compartida. Las notas aseguran que se valida contra cuda-gdb en integración continua. Suma diagnósticos de registros, memoria compartida y spills extraídos de ptxas, con un aviso cuando un kernel con límites de lanzamiento derrama registros, y un sistema de trazas de eventos dentro del kernel llamado IKET.
En el compilador hay cambios para Blackwell y Hopper: punteros de memoria compartida de 32 bits, un pipeline de siete etapas con TMA y tcgen05 en B200, y acumuladores wgmma de Hopper en registros sin spills a memoria local. NVIDIA también cifra en más de 1.240 los intrínsecos verificados, de sm_75 a sm_100a/sm_120a, y en 241 las pruebas de humo en GPU, con 16 guardas automáticas en el repositorio.
Los requisitos para cuda-oxide quedan así:
| Componente | cuda-oxide 0.3 (SIMT) | cutile-rs (Tile) |
|---|---|---|
| Rust | 1.99.1 estable, fijado, con rustc-dev y llvm-tools | Estable 1.89 o posterior |
| CUDA | Toolkit 13.0 o superior y controlador 13.x (R580+) | 13.3 recomendado |
| Compilación | Por adelantado, Rust a PTX | JIT, Rust a Tile IR y cubin |
| Otros | cargo-oxide, Clang 21 para bindgen; LLVM 21 opcional | GPU sm_80 o superior |
La documentación advierte de que el llc de LLVM 20 y anteriores no maneja los intrínsecos de TMA, tcgen05 y WGMMA que emite el compilador. Y la propia NVIDIA rebaja las expectativas: el README del repositorio lo define como un proyecto de investigación en fase temprana con fallos, funciones incompletas y cambios de API previsibles, y su página para desarrolladores lo califica de experimental.
Tampoco hay datos que muestren que el Rust compilado con cuda-oxide supere a CUDA C++. Las notas hablan de mejoras de código que, en traducción propia, buscan «cerrar la distancia con CUDA C++ ajustado a mano», lo que implica que esa distancia existe. El proyecto convive además con iniciativas comunitarias como rust-cuda, que compila Rust a NVVM IR, y Rust-GPU, orientado a SPIR-V y a gráficos; el libro de cuda-oxide afirma que NVIDIA colabora con los mantenedores de rust-cuda y presenta ambos proyectos como complementarios.
Preguntas frecuentes
¿CUDA Oxide 0.3 funciona ya con Rust estable?
Sí. Usa Rust estable 1.99.1, aunque fija esa versión y necesita los componentes rustc-dev y llvm-tools porque el backend se apoya en partes internas de rustc.
¿Qué diferencia hay entre cuda-oxide y cutile-rs?
cuda-oxide sigue el modelo SIMT y compila a PTX con control directo de hilos y memoria compartida. cutile-rs usa el modelo CUDA Tile, compila en tiempo de ejecución y deja al compilador el reparto del trabajo. Desde la 0.3 comparten runtime de host.
¿cute-rs funciona en GPU Hopper como la H100?
No. En la versión 0.3.0 solo admite sm_100a (clase B200) y sm_120a (RTX 50), y rechaza Hopper antes de compilar.
¿Está listo para producción?
NVIDIA lo describe como experimental y de fase temprana, y avisa de fallos y cambios de API. Las demostraciones prueban que funciona, no que rinda más que CUDA C++.
Fuentes:
- NVIDIA/cuda-rust: notas de la versión cuda-oxide v0.3.0 (GitHub)
- NVIDIA/cuda-rust: repositorio, README y matriz de compatibilidad de cute-rs
- Documentación oficial de CUDA Rust
- NVIDIA Developer: CUDA Rust
- NVIDIA Technical Blog: Introducing CUDA Rust: Two Tracks for Writing GPU Kernels
- Phoronix: CUDA Oxide 0.3 Vastly Improves Rust Support For NVIDIA CUDA Programming
- Rothy Yorn (Medium): NVIDIA Just Made Rust GPU Programming More Serious. CUDA Oxide 0.3 Is Why
