CUDA-oxide : le compilateur officiel Rust-to-CUDA de Nvidia

 (nvlabs.github.io)
1 points par GN⁺ 2 시간 전 | 1 commentaires | Partager sur WhatsApp
  • cuda-oxide est un compilateur expérimental qui permet d’écrire des kernels GPU SIMT en Rust idiomatique, proche de la sécurité, et de compiler directement du code Rust standard en PTX
  • Il utilise uniquement Rust, sans DSL ni bindings vers un langage externe, et suppose une compréhension de l’ownership, des traits et des génériques ; la section async requiert aussi des connaissances sur .await
  • La v0.1.0 est une version alpha initiale, il faut donc s’attendre à des bugs, des fonctionnalités incomplètes et des changements d’API cassants
  • L’exemple s’exécute avec cargo oxide run vecadd, et la fonction #[kernel] à l’intérieur de #[cuda_module] effectue une addition de vecteurs avec thread::index_1d()
  • #[cuda_module] intègre les artifacts device dans le binaire hôte, et génère un chargeur typé ainsi que des méthodes d’exécution par kernel

Mode d’utilisation et code généré

  • Démarrage rapide

    • Une fois les prérequis d’installation remplis, l’exemple se compile et s’exécute avec cargo oxide run vecadd
    • Les instructions d’installation se trouvent dans prerequisites
    • L’exemple définit la fonction #[kernel] vecadd dans un module #[cuda_module], récupère l’index avec thread::index_1d() et écrit a[i] + b[i] dans DisjointSlice<f32>
    • Côté hôte, il utilise CudaContext::new(0), le flux par défaut, kernels::load(&ctx), puis exécute le kernel avec DeviceBuffer::from_host, DeviceBuffer::<f32>::zeroed et LaunchConfig::for_num_elems(1024)
    • Le résultat est récupéré avec c.to_host_vec(&stream) et vérifié avec result[0] == 3.0

  • Fonctionnement de #[cuda_module]

    • #[cuda_module] intègre les artifacts device générés dans le binaire hôte
    • Il génère une fonction typée kernels::load et des méthodes d’exécution pour chaque kernel
    • Lorsqu’il faut charger un artifact sidecar spécifique ou créer un code d’exécution personnalisé, les API de plus bas niveau load_kernel_module et cuda_launch! restent utilisables

Prérequis et orientation

  • cuda-oxide vise à permettre l’écriture de kernels GPU à l’aide du système de types de Rust et de son modèle d’ownership, avec la sécurité comme objectif de premier plan
  • Les GPU ayant des subtilités propres, il est nécessaire de lire the safety model
  • Il ne s’agit pas d’un DSL, mais d’un backend personnalisé de génération de code rustc qui compile du Rust pur en PTX
  • Il prend en charge une exécution asynchrone en composant le travail GPU sous forme de graphe de DeviceOperation à exécution différée, planifié sur un pool de streams, puis en attendant le résultat avec .await
  • Il part du principe que vous êtes à l’aise avec l’ownership, les traits et les génériques de Rust ; les chapitres suivants sur la programmation GPU async demandent aussi des connaissances sur async/.await et des runtimes comme tokio
  • Les références proposées sont The Rust Programming Language, Rust by Example, Async Book
  • La version v0.1.0 est une release alpha initiale, et il faut s’attendre à des bugs, des fonctionnalités incomplètes et des changements d’API cassants

À lire aussi

1 commentaires

 
GN⁺ 2 시간 전
Commentaires sur Hacker News
  • C’est vraiment impressionnant. J’utilise depuis longtemps des kernels CUDA personnalisés avec https://crates.io/crates/cudarc et ça a l’air de pouvoir devenir presque un remplaçant direct
    Je suis surtout curieux de voir comment ça se comparera sur les temps de build. La plupart des crates Rust CUDA reposent sur CMake ou des appels à nvcc, ce qui peut rendre la compilation péniblement lente
    Justement, la semaine dernière, en profilant des temps de build, j’ai vu que des outils comme sccache peuvent fortement réduire les temps de recompilation via le cache des artefacts, mais le coût des appels nvcc personnalisés reste là. Par exemple, candle de Hugging Face appelle aussi une commande nvcc personnalisée pour compiler les kernels : https://arpadvoros.com/posts/2026/05/05/speeding-up-rust-whi...
    • Cudarc est vraiment bien
      Personnellement, je n’ai pas tant souffert que ça du fait que la plupart des crates Rust CUDA appellent CMake ou nvcc et compilent lentement. Si on regarde la crate cuda_setup, conçue pour gérer les scripts de build, ce n’est qu’un simple build.rs, donc ça ne recompile que lorsque les fichiers changent, et le temps de compilation reste très faible par rapport au code Rust côté CPU
    • Je me demande si d’autres ont aussi l’impression que cuda-oxide ressemble à un quasi remplaçant direct de cudarc
      Ce serait vraiment bien, mais personnellement je pense que ce sera plutôt un complément. Je me demande aussi ce qui différencie cuda-oxide, au-delà du fait qu’il soit entièrement contrôlé par NVIDIA
  • Le fait d’aller “directement vers PTX” me paraît étrange. Le NVIDIA MLIR récent est aussi plutôt bon et rapide. Ou alors ils auraient pu cibler le Tile IR [1], plus simple et très en vogue en ce moment, celui qu’utilise CuTile
    Tile IR est un peu plus haut niveau, donc bien plus facile à cibler, au prix de pertes surtout sur des points comme la fusion d’épilogue
    [1] https://docs.nvidia.com/cuda/tile-ir/
    [2] https://developer.nvidia.com/cuda/tile
  • Je suis assez curieux de voir comment le modèle mémoire de Rust a été adapté à la sémantique CUDA. J’aimerais aussi savoir ce qui change par rapport à CUDA C++, et si le système de types de Rust peut réellement apporter davantage de sûreté à CUDA
    J’ai l’impression qu’écrire des kernels GPU est intrinsèquement non sûr. Entre le fonctionnement du matériel et le besoin constant d’optimiser à l’extrême, ça semble trop difficile d’en faire un langage sûr
    • Les quatre grandes différences visibles sont les suivantes. D’abord, contrairement à l’appel manuel de cudaFree, ça gère les use-after-free et la sémantique de drop
      Ensuite, là où en C++ les arguments void* ne sont qu’un tableau de pointeurs dont on ne vérifie que le nombre, ici les arguments du kernel sont imposés via cuda_launch!
      Troisièmement, il y a le problème de l’aliasing lors des écritures mutables. En C++, un code où plusieurs threads écrivent dans out[i] avec le même i compile quand même, mais DisjointSlice et ThreadIndex n’ont pas de constructeurs publics, et on est limité aux API https://github.com/NVlabs/cuda-oxide/blob/2a03dfd9d5f3ecba52... index_1d, index_2d, index_2d_runtime
      Quatrièmement, en C++, on peut faire un cuda memcpy d’un std::string ou d’à peu près n’importe quel POD et corrompre son état, alors qu’ici seuls DisjointSlice, les scalaires et les closures sont acceptés https://nvlabs.github.io/cuda-oxide/gpu-programming/memory-a...
      Plus de détails ici : https://nvlabs.github.io/cuda-oxide/gpu-safety/the-safety-mo... et https://nvlabs.github.io/cuda-oxide/gpu-programming/memory-a.... Bien sûr, ça ne capture pas tout, mais on dirait que ça fournit bien plus de garde-fous contre le comportement indéfini que des fichiers .cu bruts
    • À noter que le modèle mémoire de Rust est presque entièrement identique à celui du C++ par conception. Les opérations atomiques sont les mêmes, et il y a aussi des notions comme la provenance
      Reste à voir si c’est un langage pratique pour programmer des GPU, mais je ne serais pas surpris qu’on puisse construire une API un peu DSL qui permette d’écrire du code sûr tout en exploitant toutes les bizarreries propres aux GPU. Au fond, CUDA n’est-il pas déjà un peu cela ?
    • C’est décrit de manière assez détaillée dans la documentation. Il y a une couche sûre, une couche en grande partie sûre et une couche non sûre
      Pour les tâches sûres mais parallèles qui s’intègrent mal dans le modèle Send/Sync de Rust, il a fallu accepter une certaine maladresse
    • J’imagine que ça dépend des objectifs. Si l’on écrit une application en Rust et qu’on veut juste faire de temps en temps du calcul GPU dedans, honnêtement, ça m’importe peu
      Si on peut profiter du modèle mémoire ou de propriété avec peu de friction, tant mieux. Mais si cela rend l’expérience d’usage beaucoup plus pénible, je n’en veux pas
      Pour moi, la référence, c’est ce que fait actuellement Cudarc. La gestion mémoire n’est pas très envahissante, et on a essentiellement une syntaxe impérative qui encapsule la FFI avec quelques lignes de script de build qui appellent nvcc quand le kernel change
  • Je me demande ce que ça signifie pour Slang[0]. L’idée de fond semble être que les gens veulent programmer les GPU avec un langage plus moderne, et là on dirait qu’il suffit désormais d’utiliser Rust
    Pour info, j’aime bien Slang
    [0]: https://shader-slang.org/
    • Écrire des shaders est, du moins pour l’instant, pratiquement différent d’écrire des kernels CUDA. Les shaders sont à la fois plus haut niveau et plus bas niveau, avec plein de particularités qui viennent du fait qu’ils visent un ensemble spécifique et limité de fonctionnalités de pilote et de GPU
      Par exemple les descriptor sets, les registres de ressources ou les limites de dispatch
    • Ce ne sont pas les mêmes cibles. Côté Slang, on s’intéresse davantage à la programmation graphique qu’aux algorithmes d’IA
      Les langages de shading sont aussi plus conviviaux sur le plan des fonctionnalités. Et puis NVIDIA utilise déjà Slang en production, donc ces équipes ne vont probablement pas réécrire leurs pipelines de shaders en Rust
  • À propos de Rust et des langages de programmation “sûrs”, je me demande si quelqu’un en sait plus sur la façon dont NVIDIA utilise Spark/Ada
    Je n’ai trouvé que ceci
    https://www.adacore.com/case-studies/nvidia-adoption-of-spar...
  • Dès la phrase “Pas de DSL, pas de bindings pour un langage étranger, juste Rust”, on a l’impression que, malgré le côté port officiel de CUDA, ils n’ont même pas vraiment soigné le paragraphe d’introduction
    J’ai quand même essayé d’ignorer ça pour lire la documentation, et au moment où ça commençait à devenir intéressant avec l’IR personnalisé, je suis tombé sur une phrase du style “une implémentation MLIR, c’est du C++ avec du TableGen, un système de build qui oblige à compiler tout LLVM, et des sessions de débogage qui vous font remettre en question vos choix de carrière”, et là j’ai eu du mal à continuer à prendre ce secteur au sérieux
    • L’ensemble de la base de code semble en grande partie écrit par IA
    • S’ils n’avaient pas utilisé l’IA pour la page web, on aurait eu droit à des réactions du style “pourquoi NVIDIA n’utilise-t-il pas l’IA pour son propre site et sa propre documentation ? Ils ne croient donc pas à leur propre discours sur les usines à IA et les milliers d’agents gérés par leurs employés ?”
      Ça ressemble exactement au bon niveau de dogfooding qu’on peut attendre d’une entreprise en plein battage autour de l’IA
    • Même le nom CUDA-oxide donne l’impression qu’ils ignorent que le nom du langage Rust vient de champignons, pas de l’oxydation
    • Je ne vois pas très bien ce qui pose problème exactement. Le fait que quelqu’un ait simplement observé que MLIR est très complexe et dépend fortement de LLVM ?
  • Des choses comme TileLang https://github.com/tile-ai/tilelang et Tile Kernels https://github.com/deepseek-ai/TileKernels finiront par rendre CUDA obsolète
    • CUDA a presque 20 ans et ne va probablement pas disparaître dans les prochaines années
    • C’est une affirmation assez énorme pour si peu d’éléments
  • Si l’on lit https://nvlabs.github.io/cuda-oxide/gpu-safety/the-safety-mo..., on voit que les kernels GPU tournent sur des milliers de threads regardant la même mémoire en même temps, et que si, côté CPU, Rust empêche les data races via la propriété et l’emprunt, sur GPU il peut y avoir 2048 threads par SM qui démarrent dans la même fonction et pointent vers le même tampon de sortie, ce pour quoi le borrow checker n’a pas été conçu
    cuda-oxide rend structurellement sûr le cas courant où “un thread écrit un élément”, et impose du unsafe avec des contrats documentés pour les cas plus rares comme la mémoire partagée, les warp shuffles ou les intrinsics matérielles, tandis que les fonctionnalités de pointe comme TMA, les tensor cores ou la communication au niveau cluster restent entièrement manuelles afin de refléter la complexité du matériel
    Mais ce n’est pas très dans l’esprit de Rust. En Rust, quand les abstractions existantes ne collent pas bien au problème, on crée de nouvelles abstractions sûres. Rust for Linux en est un exemple
    Si ce n’est pas sûr, je me demande quel est l’intérêt d’utiliser Rust. Fournir des API unsafe à ceux qui doivent arracher le dernier pour cent de performance, d’accord, mais cela ne devrait pas être le comportement par défaut
    Ça me fait penser aux bibliothèques en espace utilisateur pour des API comme io_uring ou Vulkan. Concevoir des API sûres pour ce genre de choses est assez difficile, et il y a effectivement eu des tentatives qui n’étaient pas sound
  • Je me demande si quelqu’un sait si cela permettra de partager des structures entre l’hôte et le device. C’est un gros morceau qui manquait jusqu’ici dans les workflows Rust/CUDA existants
    Idem pour la barrière de sérialisation / de bytes entre les deux
  • L’une des choses qui me rendaient prudent avec Rust dans CUDA, c’est que Rust ajoute parfois un peu de surcharge qui est généralement négligeable, mais qui peut devenir importante ici
    Par exemple, je me demande si les vérifications de limites de tableau pourraient entraîner une utilisation supplémentaire de registres et ainsi réduire la concurrence du kernel