⏱️ Lectura: 14 min
El 7 de mayo de 2026, la división de research de Nvidia, NVlabs, publicó la versión v0.1.0 de CUDA-oxide, un compilador oficial que permite escribir kernels CUDA directamente en Rust puro y compilarlos a PTX sin pasar por C++ ni por nvcc. El mismo archivo .rs contiene código CPU y código GPU, y se compila con un solo comando: cargo oxide build.
📑 En este artículo
- TL;DR
- Qué es CUDA-oxide y qué hace distinto
- Pipeline de compilación: del .rs al PTX
- La API: dos macros que cambian todo
- Lo verdaderamente novedoso: closures en GPU
- Async GPU nativo con cuda-async
- Por qué Nvidia oficializa Rust ahora
- Limitaciones reales: lo que aún no hace
- Qué sigue para CUDA-oxide
- Preguntas frecuentes
- Referencias
Es la primera vez que Nvidia respalda oficialmente Rust como lenguaje de primera clase para programación de GPUs. Hasta ahora los esfuerzos serios eran comunitarios. CUDA-oxide está en alpha experimental, pero marca un giro estratégico en el stack dominante de cómputo paralelo del planeta.
TL;DR
- Nvidia publicó CUDA-oxide v0.1.0 el 7 de mayo de 2026 bajo licencia Apache-2.0 desde su división NVlabs.
- Es un codegen backend de rustc: compila funciones marcadas con
#[kernel]directo a PTX sin pasar por C++ ninvcc. - Single-source real: el código CPU y GPU vive en el mismo archivo
.rsy se compila concargo oxide build. - Pipeline interno: Rust MIR → Pliron IR → LLVM IR → PTX. Requiere LLVM 21+ por intrinsics TMA, tcgen05 y WGMMA.
- Soporta closures con captures en GPU, algo que CUDA C++ solo logra con templates manuales y
__device__functors. - Alpha experimental: API inestable, Linux-only, nightly Rust, sin autodiff (no apto para training ML puro).
- Primera vez que Nvidia respalda oficialmente Rust en GPU; rust-cuda y cudarc eran proyectos comunitarios.
- 1.445 stars en GitHub y 148+ puntos en Hacker News al 11 de mayo, sin campaña de marketing.
Qué es CUDA-oxide y qué hace distinto
CUDA-oxide no es una biblioteca de bindings ni un DSL. Es un codegen backend personalizado para rustc: toma funciones de Rust marcadas con el macro #[kernel] y las compila directamente a PTX, el assembly intermedio de las GPUs Nvidia. No hay un sub-lenguaje aparte, no hay FFI hacia un kernel escrito en C++, y no hay paso intermedio por nvcc.
Esto lo diferencia tajantemente de los proyectos previos que intentaron acercar Rust al ecosistema CUDA:
- cudarc y rust-cuda: bindings que envuelven el runtime de CUDA. El kernel sigue escribiéndose en CUDA C++ y se carga como binario externo desde el código Rust.
- Triton (OpenAI) y Halide: DSLs embebidos. Uno escribe en un mini-lenguaje específico que se compila a GPU, no en el lenguaje host.
- ThunderKittens: librería en C++ con abstracciones de tiles para GPUs modernas, no relacionada con Rust.
CUDA-oxide es single-source en el sentido fuerte de la palabra: el código que corre en CPU y el que corre en GPU comparten archivo, funciones auxiliares, tipos y, lo más interesante, garantías del sistema de tipos de Rust. El borrow checker funciona a través del boundary host/device, no solo en cada lado por separado.
Pipeline de compilación: del .rs al PTX
El proceso de compilación de un kernel en CUDA-oxide atraviesa cinco etapas, y entender ese pipeline ayuda a anticipar dónde aparecerán bugs o limitaciones durante el ciclo alpha del proyecto.
flowchart LR
A[".rs source"] --> B["Rust MIR"]
B --> C["Pliron IR"]
C --> D["LLVM IR"]
D --> E["PTX assembly"]
E --> F["Driver Nvidia"]Cada paso cumple un rol específico:
- Rust MIR: representación intermedia de rustc, el mismo IR que usa el compilador estándar para el lado CPU. CUDA-oxide intercepta MIR en lugar de generar máquina virtual directamente.
- Pliron IR: un framework de IR extensible escrito en Rust, modelado en MLIR pero implementado nativo. Pliron permite definir dialectos, capas de IR especializadas. CUDA-oxide define dialectos propios para representar conceptos GPU como warps, bloques y memoria compartida.
- LLVM IR: el IR universal de LLVM. Aquí se aplican optimizaciones clásicas (constant folding, dead code elimination, mem2reg).
- PTX: el assembly intermedio de Nvidia. El driver de la GPU lo recompila a SASS, el código nativo real del chip, en tiempo de ejecución.
⚠️ Ojo: CUDA-oxide requiere LLVM 21 o superior. Las versiones anteriores no emiten intrinsics TMA, tcgen05 y WGMMA, necesarios para Hopper y Blackwell. Kernels simples podrían compilar con LLVM 20, pero cualquier código que toque arquitecturas modernas falla en silencio o produce PTX inválido.
El comando cargo oxide pipeline vecadd permite inspeccionar cada paso del pipeline para un kernel específico. Es la herramienta clave para entender cómo Rust se traduce a PTX cuando aparece un bug en la cadena de lowering.
La API: dos macros que cambian todo
Toda la mecánica de CUDA-oxide descansa en dos macros principales: #[cuda_module] y #[kernel]. El primero declara un módulo Rust que contiene código destinado a GPU; el segundo marca una función concreta como kernel ejecutable en device.
Cuando rustc procesa un módulo marcado con #[cuda_module], CUDA-oxide hace tres cosas: compila las funciones #[kernel] a PTX, embebe ese artefacto binario dentro del ejecutable host, y genera funciones tipadas load y métodos launch para cada kernel. Todo eso ocurre en un solo paso de cargo oxide build, sin scripts de generación intermedios ni configuración separada para device y host.
El ejemplo canónico, una suma vectorial, da la dimensión real de lo limpio que queda el código:
use cuda_device::{cuda_module, kernel, thread, DisjointSlice};
use cuda_core::{CudaContext, DeviceBuffer, LaunchConfig};
#[cuda_module]
mod kernels {
use super::*;
#[kernel]
fn vecadd(a: &[f32], b: &[f32], mut c: DisjointSlice<f32>) {
let idx = thread::index_1d();
let i = idx.get();
if let Some(c_elem) = c.get_mut(idx) {
*c_elem = a[i] + b[i];
}
}
}
fn main() {
let ctx = CudaContext::new(0).unwrap();
let stream = ctx.default_stream();
let module = kernels::load(&ctx).unwrap();
let a = DeviceBuffer::from_host(&stream, &[1.0f32; 1024]).unwrap();
let b = DeviceBuffer::from_host(&stream, &[2.0f32; 1024]).unwrap();
let mut c = DeviceBuffer::<f32>::zeroed(&stream, 1024).unwrap();
module.vecadd(&stream, LaunchConfig::for_num_elems(1024), &a, &b, &mut c).unwrap();
let result = c.to_host_vec(&stream).unwrap();
assert_eq!(result[0], 3.0);
}Vale la pena detenerse en algunos detalles. El tipo DisjointSlice es una abstracción específica de CUDA-oxide que captura, en el sistema de tipos, la invariante de que cada thread escribe en una posición distinta del slice. Esto permite que el borrow checker valide accesos concurrentes a memoria que en CUDA C++ se gestionan a mano. La función thread::index_1d() es la versión tipada del clásico blockIdx.x * blockDim.x + threadIdx.x.
El lado host queda igual de claro: CudaContext, DeviceBuffer y LaunchConfig son los equivalentes en Rust de las APIs del runtime CUDA, pero con manejo de errores idiomático mediante Result y RAII para liberar recursos automáticamente al salir de scope.
Lo verdaderamente novedoso: closures en GPU
Si tuviera que destacar la única capacidad que diferencia a CUDA-oxide de cualquier otro intento previo en cualquier lenguaje, serían las closures con captures que viajan al device. En CUDA C++, ejecutar una función arbitraria sobre cada elemento de un buffer requiere templates manuales, functors __device__ y una danza compleja con plantillas para que el compilador deduzca el tipo. En CUDA-oxide se hace así:
#[kernel]
pub fn map<T: Copy, F: Fn(T) -> T + Copy>(
f: F,
input: &[T],
mut out: DisjointSlice<T>,
) {
let idx = thread::index_1d();
let i = idx.get();
if let Some(out_elem) = out.get_mut(idx) {
*out_elem = f(input[i]);
}
}
// host side:
let factor = 2.5f32;
module.map::<f32, _>(
&stream,
LaunchConfig::for_num_elems(1024),
move |x: f32| x * factor,
&input,
&mut output,
).unwrap();La closure move |x: f32| x * factor captura la variable factor del scope host. CUDA-oxide hace dos cosas no triviales con ella: monomorphiza la closure a un tipo concreto durante la compilación, y luego scalariza los valores capturados, pasándolos como parámetros del kernel al lanzarlo. Para el desarrollador es transparente; para el runtime, es código GPU optimizado sin overhead de heap ni indirección.
💭 Clave: esto convierte a Rust en el primer lenguaje de sistemas con soporte first-class para abstracciones funcionales sobre GPU. Combinadores tipomap,filteryfoldescritos como kernels genéricos se vuelven viables sin sacrificar performance ni claridad.
El efecto secundario interesante es que las librerías de álgebra lineal o numerical computing pueden exponer APIs de alto nivel sin obligar a sus usuarios a escribir CUDA. Una closure en Rust puede expresar la operación elemento por elemento, y el compilador se encarga de moverla al device.
Async GPU nativo con cuda-async
El crate cuda-async expone un tipo DeviceOperation que representa trabajo perezoso pendiente de ejecutar en GPU. Los métodos terminados en _async devuelven una DeviceOperation que no se ejecuta hasta que se compone con otras y se invoca .sync() o .await:
let op = module.map_async::<f32, _>(
launch_cfg,
move |x: f32| x * 2.0,
&input,
&mut output,
)?;
op.sync()?; // bloqueante
// o, dentro de un runtime async:
op.await?;El detalle relevante es que DeviceOperation implementa Future, lo que permite componer trabajo GPU dentro de un runtime async como tokio sin spawn de threads dedicados ni patrones de polling manual. Múltiples kernels pueden lanzarse en paralelo sobre streams CUDA y esperarse colectivamente con join_all o select!, sin cambiar el modelo mental respecto a I/O async tradicional.
Para servidores de inferencia o pipelines de procesamiento por lotes, este modelo aproxima la programación GPU a la programación de red en Rust: perezosa, componible y sin abstracciones costosas en tiempo de ejecución. Es uno de los puntos donde el lenguaje host realmente aporta valor diferencial frente a una API equivalente en C++.
Por qué Nvidia oficializa Rust ahora
La señal estratégica detrás de CUDA-oxide es probablemente más relevante que el código mismo. Hasta esta semana, todos los esfuerzos serios de Rust sobre GPU eran impulsados por la comunidad. El proyecto original rust-cuda se detuvo en 2020. Su reinicio en 2024, bajo la organización Rust-GPU, sigue siendo trabajo voluntario sin endorso vendor. cudarc, la opción más usada en producción durante 2025, es mantenida mayoritariamente por un solo desarrollador.
Que Nvidia publique CUDA-oxide bajo licencia Apache-2.0 desde NVlabs, su división de research, y declare explícitamente que el proyecto se ofrece made available to the broader Rust community, es una declaración de intenciones. Indica que la empresa dueña del stack de cómputo paralelo dominante considera a Rust suficientemente maduro como para ser un lenguaje first-class en su ecosistema. No reemplaza a CUDA C++ — los millones de líneas existentes no se reescriben de la noche a la mañana — pero abre la puerta a que código nuevo nazca directamente en Rust.
Hay además un contexto regulatorio y de seguridad. Iniciativas como las recomendaciones de la CISA sobre lenguajes memory-safe para infraestructura crítica empujan a las grandes plataformas a ofrecer alternativas seguras a C y C++. El cómputo GPU era una frontera donde Rust no tenía respuesta oficial; ahora la tiene.
Para Nvidia, además, alinearse con Rust significa atraer a una nueva generación de desarrolladores que prefieren herramientas con garantías estáticas. La adopción de Rust en el kernel de Linux, en navegadores y en infraestructura cloud sugiere que la cohorte de talento próxima esperará tooling de ese nivel también para GPUs.
Limitaciones reales: lo que aún no hace
CUDA-oxide es honesto en su propio README respecto a su estado: expect bugs, incomplete features, API breakage as we work to improve it. Vale la pena enumerar los puntos concretos donde el proyecto todavía no compite con CUDA C++ en producción.
- Sin autodiff (diferenciación automática): este es el caveat más comentado en Hacker News. Sin AD first-class, CUDA-oxide no compite con frameworks como JAX, PyTorch o Mojo en el espacio de training de modelos de IA. Es un target apto para kernels HPC, scientific computing y procesamiento numérico, no para gradientes.
- Linux-only por ahora: el equipo testea contra Ubuntu 24.04. Otras distros podrían funcionar, pero Windows y macOS no están en el roadmap inmediato.
- Toolchain bleeding-edge: requiere Rust nightly con componentes
rust-srcyrustc-dev, LLVM 21 (recién promovido a GA), clang-21 y libclang-common-21-dev. La barrera de entrada es alta para usuarios casuales. - API inestable: v0.1.0 es alpha. Cada release minor puede romper compatibilidad hasta que el proyecto alcance 1.0.
- Sigue dependiendo del stack cerrado de Nvidia: aunque el frontend sea Rust y Apache-2.0, el driver, el runtime y el ISA siguen siendo propietarios. Para puristas open-source, CUDA-oxide cambia el frontend, no el stack completo.
📌 Nota: Una crítica recurrente en Hacker News apunta a los tiempos de compilación. Monomorphizar closures genéricas para device puede inflar el build comparado con nvcc. Aún no hay benchmarks públicos comparativos publicados.
Qué sigue para CUDA-oxide
El proyecto no publicó un roadmap formal, pero los issues abiertos y los hilos en Hacker News sugieren cuatro frentes de trabajo en los próximos meses. El primero es autodiff: la integración con Enzyme o algún backend AD propio es la petición más insistente de la comunidad, ya que abriría el camino a training de modelos en Rust puro sobre GPU Nvidia.
El segundo es estabilización de la API. Para ser viable en producción, CUDA-oxide necesita garantizar que un kernel escrito hoy seguirá compilando en seis meses. La cadencia hacia 1.0 será el indicador clave que mire la comunidad.
El tercero es rendimiento de compilación. Si los tiempos de build se vuelven prohibitivos comparados con nvcc, el proyecto quedará confinado a código nuevo y pequeño. Caching incremental de PTX y compilación paralela de kernels son áreas naturales de mejora.
El cuarto, y quizás el de mayor impacto a largo plazo, es portabilidad. Por ahora CUDA-oxide es Linux-only. La adopción real en industria requerirá soporte Windows y, eventualmente, alguna forma de target alternativo a PTX para no atarse 100% a Nvidia.
📖 Resumen en Telegram: Ver resumen
Preguntas frecuentes
¿CUDA-oxide reemplaza a nvcc para proyectos existentes?
No. CUDA-oxide está pensado para escribir kernels nuevos en Rust desde cero. Los proyectos existentes en CUDA C++ siguen necesitando nvcc. El objetivo es abrir un camino paralelo, no portar el ecosistema CUDA actual.
¿Funciona en GPUs AMD o Intel?
No. CUDA-oxide genera PTX, el assembly específico de Nvidia. Para otras GPUs habría que portar el backend a SPIR-V (Intel/Vulkan) o GCN/RDNA (AMD), tarea no contemplada por ahora.
¿Sirve para entrenar modelos de IA?
Solo parcialmente. Sin autodiff first-class, escribir entrenamiento desde cero en CUDA-oxide implica derivar gradientes manualmente. Es viable para inference y para kernels custom dentro de un framework como PyTorch, no para reemplazar a JAX o PyTorch completo.
¿Necesito conocer CUDA C++ para usar CUDA-oxide?
Ayuda, pero no es estrictamente necesario. Los conceptos centrales (kernel, thread, bloque, warp, memoria compartida) son los mismos. La diferencia está en cómo se expresan en código Rust en lugar de C++. La documentación oficial incluye un Quick Start que parte desde cero.
¿Cuándo será estable la API?
El proyecto no anunció fecha. v0.1.0 es alpha experimental con breakage esperado. Un estimado razonable basado en otros proyectos similares de Nvidia es entre 12 y 18 meses para alcanzar 1.0, pero depende del feedback de la comunidad y del ritmo de adopción.
¿Cómo se compara con Triton de OpenAI?
Triton es un DSL embebido en Python pensado para kernels de deep learning de alto nivel. CUDA-oxide es un compilador de Rust general-purpose para GPU. Triton optimiza productividad en patrones ML específicos; CUDA-oxide ofrece control fino sobre el hardware preservando las garantías del sistema de tipos de Rust.
Referencias
- Repositorio oficial CUDA-oxide en GitHub — código fuente, ejemplos y issues del proyecto bajo NVlabs.
- Documentación oficial de CUDA-oxide — Quick Start, modelo de seguridad GPU y referencia completa de la API.
- Release notes v0.1.0 — primer release público, publicado el 7 de mayo de 2026.
- Pliron — framework IR extensible inspirado en MLIR que usa CUDA-oxide como capa intermedia.
- Discusión en Hacker News — comentarios técnicos de la comunidad y análisis del proyecto.
📱 ¿Te gusta este contenido? Únete a nuestro canal de Telegram @programacion donde publicamos a diario lo más relevante de tecnología, IA y desarrollo. Resúmenes rápidos, contenido fresco todos los días.
0 Comentarios