CUDA-oxide: el compilador oficial de Rust a CUDA de Nvidia

 (nvlabs.github.io)
1 puntos por GN⁺ 2 시간 전 | 1 comentarios | Compartir por WhatsApp
  • cuda-oxide es un compilador experimental para escribir kernels GPU SIMT en Rust idiomático, cercano a la seguridad, y compilar código Rust estándar directamente a PTX
  • Usa solo Rust, sin DSL ni bindings de lenguajes externos, y asume familiaridad con ownership, traits y genéricos; la sección sobre async también requiere conocimientos de .await
  • La v0.1.0 es una versión alfa temprana, por lo que hay que esperar bugs, funciones incompletas y cambios de API incompatibles
  • El ejemplo se ejecuta con cargo oxide run vecadd, y la función #[kernel] dentro de #[cuda_module] realiza suma de vectores con thread::index_1d()
  • #[cuda_module] incluye los artefactos del dispositivo en el binario host y genera un loader con tipos y métodos de ejecución por kernel

Forma de uso y código generado

  • Inicio rápido

    • Después de cumplir con los requisitos de instalación, se puede compilar y ejecutar el ejemplo con cargo oxide run vecadd
    • La guía de instalación está en prerequisites
    • El ejemplo define la función #[kernel] vecadd dentro de un módulo #[cuda_module], obtiene el índice con thread::index_1d() y escribe a[i] + b[i] en DisjointSlice<f32>
    • Del lado del host se usan CudaContext::new(0), el stream por defecto y kernels::load(&ctx), y se ejecuta el kernel con DeviceBuffer::from_host, DeviceBuffer::<f32>::zeroed y LaunchConfig::for_num_elems(1024)
    • El resultado se recupera con c.to_host_vec(&stream) y se verifica que result[0] == 3.0

  • Cómo funciona #[cuda_module]

    • #[cuda_module] incluye en el binario host los artefactos de dispositivo generados
    • Genera la función tipada kernels::load y métodos de ejecución por kernel
    • Si se necesita cargar artefactos sidecar específicos o crear código de ejecución personalizado, también se pueden seguir usando la API de bajo nivel load_kernel_module y cuda_launch!

Requisitos y dirección del proyecto

  • cuda-oxide busca que los kernels GPU se escriban con el sistema de tipos y el modelo de ownership de Rust, y pone la seguridad como objetivo de primera clase
  • Como las GPU tienen aspectos sutiles, es necesario leer the safety model
  • No es un DSL, sino un backend personalizado de generación de código para rustc que compila Rust puro a PTX
  • Soporta ejecución asíncrona, componiendo el trabajo de GPU como un grafo de DeviceOperation con ejecución diferida, programándolo en un pool de streams y esperando los resultados con .await
  • Parte de la base de que el lector ya está familiarizado con ownership, traits y genéricos de Rust; los capítulos posteriores sobre programación GPU async también requieren conocimientos de async/.await y runtimes como tokio
  • Como material de referencia, ofrece The Rust Programming Language, Rust by Example, Async Book
  • La versión v0.1.0 está en etapa alfa temprana, así que hay que esperar bugs, funciones incompletas y cambios de API incompatibles

Lecturas relacionadas

1 comentarios

 
GN⁺ 2 시간 전
Comentarios en Hacker News
  • Esto está increíble. Llevo tiempo usando kernels CUDA personalizados y https://crates.io/crates/cudarc, y esto casi parece que podría ser un reemplazo drop-in
    Sobre todo, me da curiosidad cómo se compararán los tiempos de compilación. La mayoría de los crates de Rust para CUDA dependen de CMake o de invocar nvcc, así que compilar puede volverse dolorosamente lento
    Justo la semana pasada, perfilando tiempos de compilación, vi que herramientas como sccache pueden reducir bastante los tiempos de recompilación mediante caché de artefactos, pero aun así sigue estando el costo de las invocaciones personalizadas a nvcc. Por ejemplo, candle de Hugging Face también invoca comandos nvcc personalizados para compilar kernels: https://arpadvoros.com/posts/2026/05/05/speeding-up-rust-whi...
    • Cudarc está muy bien
      En lo personal, no he sufrido tanto eso de que la mayoría de los crates de Rust para CUDA llamen a CMake o nvcc y por eso la compilación sea lenta. Si ves el crate cuda_setup, que hice para manejar scripts de compilación, es un build.rs simple, así que solo recompila cuando cambian archivos, y el tiempo de compilación es muy pequeño comparado con el código Rust del lado CPU
    • Me pregunto si otras personas también sienten que cuda-oxide parece casi un reemplazo drop-in de cudarc
      Estaría buenísimo si fuera así, pero personalmente creo que probablemente sea más bien un complemento. También me da curiosidad cuál sería el factor diferenciador de cuda-oxide, más allá de que NVIDIA lo controle completamente
  • Me parece raro que vaya “directo a PTX”. El NVIDIA MLIR reciente también es bastante bueno y rápido. O podrían haber apuntado al Tile IR [1], que es más fácil y más moderno, y que usa CuTile
    Tile IR está un poco más arriba en nivel de abstracción, así que es mucho más fácil apuntarle, y solo sales perdiendo en cosas como la fusión de epílogos
    [1] https://docs.nvidia.com/cuda/tile-ir/
    [2] https://developer.nvidia.com/cuda/tile
  • Tengo bastante curiosidad por cómo adaptaron el modelo de memoria de Rust a la semántica de CUDA. Quisiera saber qué cambia frente a CUDA C++ y si el sistema de tipos de Rust realmente puede aportar más seguridad a CUDA
    Me parece que escribir kernels para GPU es intrínsecamente inseguro. Por la forma en que funciona el hardware y porque siempre hay que optimizar al límite, suena dificilísimo hacer un lenguaje seguro para esto
    • A grandes rasgos, veo 4 diferencias. Primero, maneja use-after-free y la semántica de drop, a diferencia de llamar manualmente a cudaFree
      Segundo, mientras que en C++ los argumentos void* son un arreglo de punteros y solo se valida la cantidad, aquí cuda_launch! fuerza los argumentos del kernel
      Tercero, el problema del aliasing en escrituras mutables. En C++, incluso compila código donde dos o más hilos escriben a out[i] con el mismo i, pero DisjointSlice y ThreadIndex no tienen constructores públicos y solo permiten usar las APIs https://github.com/NVlabs/cuda-oxide/blob/2a03dfd9d5f3ecba52... index_1d, index_2d, index_2d_runtime
      Cuarto, en C++ puedes hacer cuda memcpy de std::string o prácticamente cualquier POD y corromper el estado, pero aquí solo acepta DisjointSlice, escalares y closures https://nvlabs.github.io/cuda-oxide/gpu-programming/memory-a...
      Hay más detalles en https://nvlabs.github.io/cuda-oxide/gpu-safety/the-safety-mo... y https://nvlabs.github.io/cuda-oxide/gpu-programming/memory-a.... Claro, no detecta todo, pero sí parece dar muchas más barandillas de seguridad contra comportamiento indefinido que archivos .cu raw
    • Como referencia, el modelo de memoria de Rust es intencionalmente casi idéntico al de C++. Las operaciones atómicas también son iguales, y también existen conceptos como la proveniencia
      Falta ver si es un lenguaje cómodo para programación GPU, pero no me sorprendería que se pudiera construir una especie de DSL o API razonable para escribir código seguro mientras se aprovechan todas las rarezas específicas de GPU. Al final, CUDA también es algo así, ¿no?
    • Está bastante detallado en la documentación. Hay una capa segura, una capa mayormente segura y una capa insegura
      Hace falta un poco de tosquedad para los trabajos seguros pero paralelos que no encajan fácilmente en el modelo Send/Sync de Rust
    • Supongo que depende de cuál sea tu objetivo. Si quieres escribir una aplicación en Rust y usar de vez en cuando cómputo GPU dentro de ella, la verdad no me importa demasiado
      Estaría bien poder aprovechar el modelo de memoria o de ownership con poca fricción. Pero si eso vuelve muy incómoda la experiencia de uso, no lo querría así
      Para mí, la línea base es lo que hace hoy Cudarc: no se mete demasiado con la gestión de memoria, y básicamente ofrece una sintaxis imperativa que envuelve FFI y unas cuantas líneas de script de compilación que invocan nvcc cuando cambia el kernel
  • Me pregunto qué significa esto para Slang[0]. Parecería que la idea es que la gente quiere programar GPU con un lenguaje más moderno, pero ahora se ve como que simplemente podrías usar Rust
    Dicho eso, Slang me gusta bastante
    [0]: https://shader-slang.org/
    • Escribir shaders es, al menos por ahora, sustancialmente distinto de escribir kernels CUDA. Los shaders son a la vez de más alto nivel y de más bajo nivel, y tienen muchas rarezas como resultado de estar diseñados para un conjunto específico y limitado de capacidades de driver/GPU
      Por ejemplo, descriptor sets, registros de recursos, límites de dispatch y cosas así
    • Apuntan a cosas distintas. La gente de Slang está más interesada en programación gráfica que en algoritmos de IA
      Los lenguajes de shading también son más amigables para el usuario en términos de capacidades. Además, NVIDIA ya usa Slang en producción, y esa gente no va a reescribir su pipeline de shaders en Rust
  • Relacionado con Rust y los lenguajes de programación “seguros”, me pregunto si alguien sabe más sobre cómo NVIDIA usa Spark/Ada
    Lo único que pude encontrar es esto
    https://www.adacore.com/case-studies/nvidia-adoption-of-spar...
  • Desde la frase “sin DSL, sin bindings de lenguaje extranjero, solo Rust”, ya da la impresión de que, aunque se supone que es un port oficial de CUDA, ni siquiera le pusieron atención al párrafo de introducción
    Igual intenté ignorarlo y leer la documentación, y justo cuando el IR personalizado empezaba a ponerse interesante, me topé con una línea del estilo “la implementación de MLIR fue C++ con un toque de TableGen, un sistema de compilación que te obliga a compilar todo LLVM y sesiones de depuración que te hacen cuestionar tus decisiones de carrera”, y desde ahí ya me cuesta tomarme en serio esta industria
    • Todo el codebase parece haber sido escrito con IA en gran parte
    • Si no hubieran usado IA en la web, la reacción habría sido “¿por qué NVIDIA no usa IA para escribir su propio sitio y su documentación? ¿No cree en su propia historia de empleados gestionando fábricas de IA y miles de agentes?”
      Esto se ve como el uso exacto de sus propios productos que uno esperaría de una empresa subida al hype de la IA
    • Además le pusieron CUDA-oxide, lo que delata que no saben que el nombre Rust no viene de la oxidación sino de un hongo
    • No entiendo exactamente cuál es la molestia. ¿Que alguien observó que MLIR es muy complejo y depende fuertemente de LLVM?
  • Cosas como TileLang https://github.com/tile-ai/tilelang y Tile Kernels https://github.com/deepseek-ai/TileKernels en algún momento van a volver obsoleto a CUDA
    • CUDA ya tiene casi 20 años y no va a desaparecer en varios años más
    • Es una afirmación bastante grande con muy poca evidencia
  • Viendo el documento https://nvlabs.github.io/cuda-oxide/gpu-safety/the-safety-mo..., dicen que los kernels de GPU corren en miles de hilos viendo la misma memoria al mismo tiempo, y que mientras en CPU Rust evita carreras de datos con ownership y préstamos, en GPU hay 2048 hilos por SM arrancando en la misma función y apuntando al mismo búfer de salida, así que el borrow checker no fue diseñado para esto
    cuda-oxide vuelve estructuralmente seguro el caso común de “un hilo escribe un elemento”, y para los casos menos comunes como memoria compartida, warp shuffles y funciones intrínsecas de hardware, exige unsafe con contratos documentados; mientras que para funciones de frontera como TMA, tensor cores y comunicación a nivel de clúster, lo deja totalmente manual, en línea con la complejidad del hardware
    Pero esto no se siente muy a lo Rust. En Rust, cuando las abstracciones existentes no encajan bien con el problema, se crean nuevas abstracciones seguras. Rust for Linux es un ejemplo
    Si no va a ser seguro, entonces me pregunto cuál es el punto de usar Rust. Está bien ofrecer APIs unsafe para la gente que necesita exprimir el último poco de rendimiento, pero eso no debería ser el valor por defecto
    Me hace pensar en librerías de espacio de usuario para APIs como io_uring o Vulkan. Diseñar APIs seguras para esas cosas es bastante difícil, y de hecho ha habido intentos que ni siquiera resultaron sound
  • Me pregunto si esto permitirá compartir structs entre host y device. Esa ha sido una gran pieza faltante en los flujos de trabajo Rust/CUDA hasta ahora
    Igual con la barrera de serialización/bytes entre ambos
  • La parte que siempre me ha dado cautela al usar Rust con CUDA es que Rust añade un poco de overhead que normalmente puedes ignorar, pero que aquí sí podría importar
    Por ejemplo, me pregunto si las comprobaciones de límites de arreglos podrían provocar uso extra de registros y reducir la concurrencia del kernel