CUDA Rust : Nvidia ouvre la programmation GPU aux Rustacéens

La programmation GPU a longtemps ressemblé à un club privé : bienvenue si tu parles C, CUDA, et que tu n’as pas peur des segfaults à 3h du matin. Nvidia vient de changer la donne en annonçant un support natif de Rust pour écrire des kernels GPU — ce qu’ils appellent CUDA Rust. Pas un binding hasardeux autour d’une API C, pas un wrapper communautaire sans garantie de support : du vrai Rust compilé vers du code GPU, officiel et maintenu par Nvidia.

L’annonce a buzzé sur Hacker News avec 662 points et 253 commentaires — c’est significatif pour un sujet aussi technique. La communauté Rust avait déjà ses projets maisons comme rust-cuda ou cudarc, mais là on bascule dans une autre catégorie : Nvidia met ses propres ingénieurs dessus. La différence entre un projet communautaire enthousiaste et un support éditeur, c’est souvent la question de la durée et de la maintenance.

Pour comprendre pourquoi c’est notable, rappelle-toi que le GPU reste un des derniers bastions du code vraiment dangereux : accès mémoire concurrent sur des dizaines de milliers de threads, gestion manuelle de la mémoire partagée, race conditions qui produisent des résultats silencieusement faux sans le moindre crash. Rust, avec son borrow checker et son système de types, a précisément été conçu pour éliminer ces classes d’erreurs. La question était : est-ce que ces garanties survivent au modèle de programmation GPU ?

Deux tracks, une vraie philosophie

L’annonce de Nvidia présente deux approches distinctes pour écrire des kernels en Rust — et ce choix architectural dit beaucoup sur la maturité de leur démarche.

Le premier track, « safe Rust », propose des abstractions haut niveau qui garantissent la sécurité mémoire même en contexte GPU. Tu écris du Rust normal, le compilateur te protège des erreurs classiques, et les kernels restent lisibles. C’est l’entrée dans GPU programming pour un dev Rust qui n’a jamais mis les mains dans CUDA — et c’est probablement là que 80 % des utilisateurs resteront.

Le second track, « raw CUDA Rust », te donne accès direct aux primitives CUDA — shared memory, synchronisation de threads, intrinsics bas niveau. Tu perds une partie des garanties de sécurité en échange du contrôle total. C’est là que tu vas si tu dois extraire chaque nanoseconde d’une passe de matrix multiplication ou d’un kernel de convolution pour un réseau de neurones en production.

Ce double track n’est pas un compromis bancal : c’est exactement le même modèle que Rust utilise pour la programmation système ou l’embarqué. Tu commences safe, tu passes en unsafe uniquement quand tu sais ce que tu fais, et tu encapsules ça proprement derrière une interface sure. Appliquer ce paradigme au GPU est logique — même si l’implémentation est plus complexe, parce que la hiérarchie mémoire GPU (registres, shared memory, L2, global memory) impose des contraintes très différentes du CPU.

Le borrow checker face aux threads CUDA

La vraie question technique : est-ce que les garanties de Rust ont du sens sur un GPU ? Un kernel CUDA lance potentiellement des dizaines de milliers de threads en parallèle. Le borrow checker de Rust, conçu pour empêcher les data races entre threads CPU, a-t-il quelque chose à dire sur cette architecture massivement parallèle ?

La réponse courte : oui, mais avec des nuances. Les accès à la global memory GPU peuvent bénéficier des garanties Rust — si le compilateur peut prouver statiquement qu’un seul thread écrit à une adresse donnée, il élimine une catégorie entière de bugs. En revanche, la shared memory — cette petite mémoire ultra-rapide partagée entre les threads d’un même bloc CUDA — nécessite des synchronisations explicites (__syncthreads()), et là le borrow checker seul ne peut pas tout résoudre.

C’est pourquoi le safe track introduit probablement des types qui encodent ces contraintes — un peu comme Arc<Mutex<T>> encode le partage sécurisé entre threads CPU. La shared memory devient un type spécial dont les accès sont typés-safe dans les limites de ce que le compilateur peut vérifier statiquement. C’est élégant sur le papier ; la mise en œuvre réelle sera l’épreuve de vérité une fois les projets complexes débarqués.

La vraie valeur de CUDA Rust n’est pas d’éliminer tous les bugs GPU — c’est de remonter plus tôt les erreurs qui aujourd’hui te donnent des résultats silencieusement incorrects sans le moindre signal d’alarme.

Ce que ça change pour l’écosystème IA et HPC

Le timing n’est pas anodin. Le HPC et l’IA sont aujourd’hui les deux domaines qui consomment le plus de cycles GPU — et tous les deux souffrent du même problème : le code CUDA C++ est puissant mais coûteux à maintenir, difficile à débugger, et pénible à intégrer dans des chaînes de build modernes. La friction est réelle et quotidienne pour les équipes qui l’utilisent.

Rust est déjà présent dans l’écosystème IA, notamment via des projets comme Candle (la lib d’inférence de Hugging Face) ou Burn, un framework de deep learning entièrement en Rust. Ces projets se heurtaient jusqu’ici à la barrière des kernels GPU : obligés de passer par des bindings C++ ou des compute shaders WebGPU, avec les limitations que ça implique en termes de contrôle sur les performances. CUDA Rust pourrait lever cet obstacle proprement.

Pour le HPC scientifique, c’est aussi une ouverture concrète. Beaucoup de codes de simulation sont écrits en C++ vieillissant, avec des patterns d’accès mémoire dangereux qui ont survécu faute d’alternative crédible. La possibilité d’écrire des kernels GPU en Rust — et de les intégrer progressivement via l’interop Rust/C dans une base de code existante — ouvre une voie de migration réaliste. Si tu travailles sur une application sur mesure avec du traitement intensif de données, c’est un changement d’horizon potentiellement important.

À quoi ressemble un kernel en Rust, concrètement ?

L’API exacte sera précisée au fil des releases, mais l’esprit du safe track ressemble probablement à quelque chose dans ce genre pour une simple addition de vecteurs :

#[kernel]
pub fn vector_add(a: &[f32], b: &[f32], c: &mut [f32]) {
    let idx = thread::index_1d();
    if idx < a.len() {
        c[idx] = a[idx] + b[idx];
    }
}

Comparé à l’équivalent CUDA C++ avec ses __global__, ses blockIdx.x * blockDim.x + threadIdx.x et sa gestion manuelle des dimensions de grille — c’est du jour au nuit en lisibilité. Le compilateur infère la configuration de lancement, vérifie les bounds, et génère le PTX correspondant. Le kernel devient lisible par un dev qui n’a jamais touché CUDA.

Le track raw, lui, t’expose les primitives CUDA directement avec une syntaxe Rust mais sans les protections automatiques. C’est là que tu ajustes le tile size, que tu joues avec les warp shuffles, que tu optimises l’accès coalescé à la global memory. Moins lisible, mais potentiellement aussi rapide que du CUDA C++ bien écrit. C’est le deal habituel avec Rust : tu peux aller aussi bas que le C, mais tu le sais quand tu le fais — et tu l’assumes explicitement.

Ce qui manque encore — soyons honnêtes

CUDA Rust en est à ses débuts, et ça se verra dans les premières semaines d’usage. Le tooling reste à construire : les profilers comme Nsight devront comprendre le code Rust compilé, les debuggers devront afficher correctement les stacks Rust sur device, et l’intégration avec les grandes librairies CUDA (cuBLAS, cuDNN, NCCL) reste floue. Ce n’est pas rien — ce sont les outils que les équipes utilisent tous les jours.

La chaîne de compilation soulève aussi des questions concrètes. CUDA Rust passe aujourd’hui par rustc combiné à une toolchain Nvidia spécifique. Est-ce que ça rentrera dans Cargo de façon transparente ? Les targets seront-elles supportées proprement en cross-compilation dans des environnements CI ? Ces questions détermineront l’adoption réelle bien plus que les capacités théoriques du langage.

Il y a enfin la question de la portabilité, fondamentale pour certains projets. CUDA est Nvidia-only, sans exception. Les alternatives comme wgpu (portable sur toutes les plateformes via WebGPU) ou ROCm pour AMD gardent toute leur pertinence. CUDA Rust s’adresse aux projets qui ont déjà fait le choix — ou qui ont les moyens de faire le choix — de l’écosystème Nvidia vert.

Une ligne franchie, pas une fin

CUDA Rust est une de ces annonces qui semblent anecdotiques à court terme et fondamentales à moyen terme. On n’est pas à la veille de voir les équipes ML migrer leurs kernels PyTorch de C++ vers Rust. Mais on vient de franchir une ligne réelle : Nvidia reconnaît officiellement Rust comme langage de première classe pour GPU. C’est une validation qui compte dans un écosystème qui pèse des milliards de dollars d’infrastructure.

Ce qui me convainc dans cette approche, c’est précisément le modèle à deux tracks. Il montre une maturité de conception : ne pas promettre des garanties impossibles, mais offrir une graduation entre sécurité et contrôle. C’est la bonne approche pour un domaine où les performances brutes sont non-négociables — et ça ressemble beaucoup à la philosophie Rust elle-même, ce qui est cohérent avec l’écosystème visé.

Mon pari : d’ici 18 à 24 mois, les projets Rust IA comme Candle ou Burn auront des backends CUDA Rust natifs, avec des kernels personnalisés comme vrai avantage compétitif — la capacité d’écrire des opérations spécialisées pour un modèle sans jamais toucher à du C++. L’annonce d’aujourd’hui est le premier caillou. La question ouverte, c’est de savoir si Nvidia tiendra la cadence sur le tooling. Parce qu’un langage sans debugger fonctionnel, c’est juste un langage de plus qu’on abandonne au premier bug incompréhensible.