NVIDIA AI acaba de lanzar cuda-oxide: un backend experimental del compilador Rust-to-CUDA que compila núcleos de GPU SIMT directamente en PTX

Configurar la cadena de herramientas nocturnas de Rust

cuda-oxide requiere Rust todas las noches con dos componentes adicionales: Rust-src y Rustc-dev. La cadena de herramientas está fijada en nightly-2026-04-03 a través de rust-toolchain.toml en el repositorio; se instalará automáticamente cuando ejecute por primera vez una compilación dentro del repositorio.

Si necesita instalarlo manualmente:

# Instalar la cadena de herramientas nocturna fijada
Instalación de la cadena de herramientas Rustup todas las noches-2026-04-03

# Agregar componentes requeridos
componente oxidado agregar rust-src rustc-dev \ –toolchain nightly-2026-04-03

# Confirmar que la cadena de herramientas está activa
espectáculo oxidado

ⓘ ¿Por qué estos componentes? Rustc-dev expone las API del compilador interno a las que se conecta el backend de codegen personalizado. Rust-src es necesario para que el compilador pueda encontrar y compilar sus propias fuentes de biblioteca estándar para el dispositivo de destino.

Instale LLVM 21 con el backend NVPTX

La canalización de cuda-óxido emite LLVM IR textuales (archivos .ll) y los entrega al binario llc externo para producir PTX. Necesita LLVM 21 o posterior con el backend NVPTX habilitado.

#Ubuntu/Debian
sudo apto instalar llvm-21

# Verificar que el backend NVPTX esté presente
llc-21 –versión | grep nvptx

La canalización descubre automáticamente llc-22 y llc-21 en su RUTA en ese orden. Para anclar un binario específico, configure la variable de entorno:

# Anclar a un binario de llc específico
exportar CUDA_OXIDE_LLC=/usr/bin/llc-21

⚠ Fallo común Si NVPTX no aparece en el resultado de llc-21 –version, su compilación LLVM se compiló sin el destino NVPTX. Instálelo desde el repositorio oficial de LLVM apt en lugar de los paquetes predeterminados de su distribución, que pueden omitir los backends de GPU.

Instale Clang 21 para la caja de enlaces cuda

La caja de enlaces cuda utiliza bindgen para generar enlaces FFI a cuda.h en el momento de la compilación. bindgen necesita libclang y, específicamente, necesita el propio directorio de recursos de Clang (que incluye stddef.h). Un paquete de tiempo de ejecución libclang1-* no es suficiente.

# Instale el paquete clang-21 completo (incluye encabezados de recursos)
sudo apto instalar clang-21

# Alternativamente, el paquete de encabezado -dev también funciona
sudo apto instalar libclang-common-21-dev

⚠ Síntoma de falta de sonido metálico Si solo instala el tiempo de ejecución pero no los encabezados, la compilación del host fallará con un error críptico de archivo ‘stddef.h’ no encontrado durante bindgen. Ejecute cargo oxide doctor en el siguiente paso para detectar esto antes de intentar una construcción.

Clonar el repositorio e instalar cargo-óxido

cargo-oxide es un subcomando Cargo que impulsa todo el proceso de construcción: ejecuta la construcción de óxido de carga, la ejecución de óxido de carga, la depuración de óxido de carga y el proceso de óxido de carga.

Dentro del repositorio (para probar ejemplos):

clon de git https://github.com/NVlabs/cuda-oxide.git cd cuda-óxido

# el óxido de carga funciona de inmediato a través de un alias de espacio de trabajo
ejecución de óxido de carga vecadd

Fuera del repositorio (para sus propios proyectos):

# Instalar globalmente desde la fuente git
instalación de carga \ –git https://github.com/NVlabs/cuda-oxide.git \ cargo-oxide

# En la primera ejecución, cargo-oxide busca y construye el backend de codegen

Luego verifique que se cumplan todos los requisitos previos con la verificación de estado incorporada:

médico de óxido de carga

ⓘ Lo que verifica el médico Valida su cadena de herramientas Rust (nightly, Rust-src, Rustc-dev), el kit de herramientas CUDA, la versión LLVM y la compatibilidad con NVPTX, los encabezados Clang/libclang y el binario backend de Codegen. Arregle los elementos rojos antes de continuar.

Construya y ejecute el ejemplo vecadd

El primer ejemplo canónico es vecadd, un núcleo de suma de vectores que agrega dos matrices de 1024 valores f32 en la GPU y verifica el resultado en el host.

# Construir y ejecutar de un extremo a otro
ejecución de óxido de carga vecadd

Si todo está configurado correctamente, verá:

✓ ÉXITO: ​​¡Los 1024 elementos correctos!

Para ver el proceso de compilación completo, desde Rust MIR pasando por cada dialecto de Pliron hasta PTX, ejecute:

# Imprima el rastro completo de Rust MIR – dialect-mir – mem2reg – dialect-llvm – LLVM IR – PTX
tubería de óxido de carga vecadd

Para depurar con cuda-gdb:

carga óxido depuración vecadd –tui

ⓘ Artefactos de salida Una compilación exitosa produce dos archivos: target/debug/vecadd (el binario del host) y target/debug/vecadd.ptx (el código del dispositivo). El binario del host carga el archivo PTX a través del controlador CUDA en tiempo de ejecución.

Escribiendo tu propio #[kernel] Función

Una función del núcleo está anotada con #[kernel]. Utilice DisjointSlice para salidas mutables y &[T] para entradas de sólo lectura. Acceda al índice de hardware único del hilo con thread::index_1d().

usar dispositivo_cuda::{núcleohilo, DisjointSlice};

// Seguridad de nivel 1: libre de carreras por construcción, no se necesita “inseguridad”. // DisjointSlice::get_mut() solo acepta un ThreadIndex — // un tipo opaco derivado del hardware que garantiza escrituras únicas por subproceso.
#[kernel]
pub escala(aporte: &[f32]factor: f32, mut afuera: Rebanada disjunta<f32>) {
dejar idx = hilo::índice_1d();
si lo dejamos Alguno(elem) = fuera.get_mut(idx) { *elem = entrada[idx.get()] * factor; } }

ⓘ Seguridad de nivel 1: cómo funciona ThreadIndex es un nuevo tipo opaco sobre uso que solo se puede crear a partir de registros integrados en el hardware (threadIdx, blockIdx, blockDim). Dado que cada hilo obtiene un valor único, y DisjointSlice::get_mut() solo acepta un ThreadIndex, las escrituras son libres de razas por construcción, no son inseguras en ninguna parte del núcleo.

Lanzar el kernel desde el código del host

El código del host y del dispositivo se encuentra en el mismo archivo .rs. ¡El lado del host usa CudaContext, DeviceBuffer y cuda_launch! Macro para administrar la memoria y el envío de la GPU.

usar cuda_core::{CudaContext, DeviceBuffer, LaunchConfig};
usar cuda_host::{cuda_launch, load_kernel_module};

fn principal() {
// Inicializa el contexto de GPU en el dispositivo 0
dejar ctx = CudaContexto::nuevo(0).desenvolver();
dejar corriente = ctx.flujo_predeterminado();
dejar módulo = módulo_kernel_carga(& ctx, “ejemplo_escala”).desenvolver();

// Carga datos de entrada a la memoria de la GPU
dejar datos: vec<f32> = (0..1024).mapa(|yo| yo como f32).recolectar();
dejar entrada = DispositivoBuffer::desde_host(&transmisión, &datos).desenvolver();
deja que mute salida = DispositivoBuffer::<f32>::puesto a cero(&arroyo, 1024).desenvolver();

// Enviamos el kernel: LaunchConfig ajusta automáticamente el tamaño de los bloques/cuadrículas
cuda_launch! { kernel: escala, flujo: flujo, módulo: módulo, configuración: Configuración de lanzamiento::para_num_elems(1024), argumentos: [slice(input), 2.5f32, slice_mut(output)]
}.desenvolver();

// Descargar el resultado al host
dejar resultado = salida.to_host_vec(&arroyo).desenvolver(); afirmar!((resultado[1] – 2.5).abdominales() < 1e-5); imprimir!(” ✓ ¡El kernel se ejecutó correctamente!”); }

ⓘ ¡Qué cuda_launch! escalariza la lista de argumentos (aplanamiento de sectores, escalares y cierres capturados) en parámetros del kernel PTX y distribuye el kernel en la secuencia dada. No se requiere organización manual de argumentos.

Qué explorar a continuación

Tienes una configuración de óxido de cuda que funciona. Estos son los caminos de alto valor a seguir, ordenados por complejidad:

Núcleos genéricos con monomorfización: pruebe el ejemplo genérico (ejecución genérica de óxido de carga) para ver cómo se compila fn scale para separar los núcleos PTX por tipo. Cierres con capturas: el ejemplo host_closure muestra cómo un movimiento |x: f32| El cierre del factor x * se escalariza y se pasa automáticamente como parámetros del kernel PTX. Ejecución asíncrona de GPU: cuda_launch_async! devuelve una DeviceOperation diferida que se ejecuta en .sync() o .await. Vea los ejemplos de async_mlp y async_vecadd. La memoria compartida y los intrínsecos de deformación requieren bloques inseguros con contratos de seguridad documentados. Consulte el Nivel 2 en la documentación del modelo de seguridad. GEMM a la velocidad de la luz: el ejemplo de gemm_sol logra 868 TFLOPS en B200 (58 % de cuBLAS SoL) usando cta_group::2, CLC y una canalización de 4 etapas. Núcleos tensoriales de Blackwell: el ejemplo tcgen05 apunta a sm_100a con TMEM, MMA y cta_group::2. Requiere LLVM 21+.

ⓘ La limitación conocida en v0.1.0 index_2d(stride) está documentada como incorrecta actualmente: si los subprocesos en el mismo núcleo usan diferentes valores de zancada, dos subprocesos pueden llevar &mut T al mismo elemento sin que haya peligro a la vista. Hasta que llegue la solución (levantando la zancada a un parámetro de tipo), vincule la zancada a un enlace let único y reutilícelo en cada sitio de llamada.

Documentación completa: nvlabs.github.io/cuda-oxide · Fuente: github.com/NVlabs/cuda-oxide