CUDA Oxide 0.3 passe à Rust stable et réunit SIMT et Tile dans un même runtime

J. Actualclo

X (Twitter) Facebook Pinterest LinkedIn Email

NVIDIA a publié le 31 octobre 2026 la version 0.3.0 de cuda-oxide, le compilateur qui permet d’écrire des kernels CUDA en Rust, et abandonne ainsi la dépendance aux versions nightly du langage : le projet se compile désormais avec Rust stable 1.99.1 et partage un même runtime d’hôte avec cutile-rs au sein du monorepo NVIDIA/cuda-rust.

Les points clés de CUDA Oxide 0.3 en 30 secondes

NVIDIA a publié CUDA Oxide 0.3 le 31/10/2026, avec près de 600 commits de plus de 45 contributeurs depuis la 0.2.1.
Le compilateur SIMT abandonne les versions nightly et fixe Rust stable 1.99.1, même s’il nécessite toujours rustc-dev, les outils LLVM et CUDA 13.0 ou supérieur.
cuda-oxide et cutile-rs partagent désormais un runtime d’hôte avec buffers, streams, événements et CUDA graphs.
cute-rs, encore en préversion, apporte CuTe à Rust uniquement pour Blackwell sm_100a et sm_120a.
NVIDIA maintient le projet comme expérimental.

Selon les notes de version sur GitHub, il s’agit de la plus importante livraison du projet à ce jour. NVIDIA a présenté CUDA Rust en septembre comme une plateforme offrant deux façons d’écrire des kernels, et nous avions alors raconté comment NVIDIA apportait CUDA à Rust sans quitter le langage. Cette première étape exigeait une toolchain nightly précise pour cuda-oxide. La 0.3 change cette base, sans pour autant la simplifier complètement.

Le compilateur continue de fixer sa version de Rust dans un fichier rust-toolchain.toml et nécessite les composants rustc-dev, rust-src et llvm-tools. La raison est technique : cuda-oxide est un backend de génération de code pour rustc et s’appuie sur des composants internes du compilateur (la bibliothèque librustc_driver et rustc_public), qui n’offrent aucune garantie de stabilité d’une version à l’autre. L’outil cargo-oxide gère cette toolchain fixée. La différence pratique est qu’il n’est plus nécessaire de suivre un nightly qui change tous les jours.

SIMT et Tile : deux compilateurs sur un runtime commun

Le dépôt NVIDIA/cuda-rust regroupe deux modèles de programmation. cuda-oxide suit le modèle SIMT (Single Instruction, Multiple Threads, une instruction exécutée par de nombreux threads), le modèle classique de CUDA C++ : le développeur écrit ce que fait un thread et contrôle les blocs, les warps et la mémoire partagée. Il compile en amont (ahead-of-time) de Rust vers PTX (Parallel Thread Execution, le langage intermédiaire des GPU NVIDIA) en passant par la représentation MIR de rustc, le framework d’IR Pliron et LLVM.

cutile-rs implémente le modèle CUDA Tile, que NVIDIA propose déjà en C++ et en Python. Ici, le kernel travaille sur des blocs de données entiers (tiles) et le compilateur décide de leur répartition selon chaque architecture ; le code est compilé au premier lancement vers Tile IR et un binaire cubin. Nous avions expliqué en son temps le fonctionnement de CUDA Tile. cutile-rs n’a pas besoin de la toolchain fixée : il fonctionne avec Rust stable 1.89 ou supérieur, et sa documentation recommande CUDA 13.3, indispensable pour le packaging FP4 et les MMA à mise à l’échelle par blocs.

La nouveauté de la 0.3 est que les deux partagent le runtime d’hôte formé de cuda-bindings, cuda-core, cuda-core-derive et cuda-async (version 0.5.0). Les notes de version décrivent un seul DeviceBuffer, Stream, Event, CudaGraph et DeviceOp asynchrone pour les deux modèles, de sorte qu’un kernel Tile et un kernel SIMT peuvent travailler dans le même stream sur les mêmes buffers. Elle ajoute également la gestion de la mémoire virtuelle, l’accès entre GPU, le multicast via NVLink Switch et les priorités de stream, ainsi qu’un attribut #[tile_fn] pour exporter des fonctions SIMT sous forme de bibliothèques qu’un kernel Tile peut appeler.

cute-rs : CUTLASS en Rust, uniquement pour deux variantes de Blackwell

L’élément le plus marquant est cute-rs-preview, une implémentation expérimentale de CuTe, la couche de layouts de CUTLASS (la bibliothèque de templates de NVIDIA pour l’algèbre linéaire sur GPU). Elle compile en amont grâce au compilateur fixé de CUTLASS 4.7 et couvre, selon sa propre matrice de compatibilité, un sous-ensemble validé de CuTeDSL 4.7, sans parité de fonctionnalités.

Les cibles sont strictes. Dans la 0.3.0, elle ne prend en charge que sm_100a (Blackwell pour centres de données, classe B200) et sm_120a (GeForce RTX 50). Ampere, Ada et Hopper (sm_90a) sont rejetés avant la compilation, tout comme d’autres variantes de Blackwell. La matrice indique en outre l’absence de prise en charge de wgmma, l’instruction de Tensor Cores propre à Hopper, ainsi que des kernels groupés, du MoE, des entrées mixtes ou du FP8. Sur la branche principale, après la publication, sm_103a (classe B300) figure déjà également, mais ne fait pas partie de la version 0.3.0.

Ce qu’elle offre, en revanche, c’est la possibilité d’exprimer dans le système de types de Rust les layouts, les descripteurs TMA (Tensor Memory Accelerator, le moteur de copies asynchrones de mémoire), les pipelines producteur-consommateur et les MMA des Tensor Cores, et de les vérifier avant de générer le code. La version inclut 14 kernels d’exemple et deux modèles complets qui fonctionnent sur une RTX 5090 : Qwen3-8B de 36 couches en BF16 (bfloat16, virgule flottante 16 bits) et NVFP4 (le format 4 bits à mise à l’échelle par blocs de NVIDIA), et Qwen3-TTS 1.7B. En NVFP4, les poids occupent 3,64 GiB contre 12,94 GiB en BF16, soit 3,56 fois moins.

Débogage avec cuda-gdb, tests et prérequis d’installation

Pour ceux qui doivent maintenir des kernels, la 0.3 apporte un débogage DWARF (le format standard d’informations de débogage) complet avec cuda-gdb via les options --device-debug et --lineinfo, qui conservent les emplacements de code, les closures, les variantes d’enum et la mémoire partagée. Les notes affirment que ce débogage est validé contre cuda-gdb en intégration continue. Elle ajoute des diagnostics de registres, de mémoire partagée et de spills extraits de ptxas, avec un avertissement lorsqu’un kernel avec des limites de lancement déborde des registres, ainsi qu’un système de traces d’événements dans le kernel nommé IKET.

Côté compilateur, des changements concernent Blackwell et Hopper : pointeurs de mémoire partagée 32 bits, pipeline à sept étapes avec TMA et tcgen05 sur B200, et accumulateurs wgmma de Hopper dans des registres sans spills vers la mémoire locale. NVIDIA chiffre aussi à plus de 1 240 les intrinsèques vérifiés, de sm_75 à sm_100a/sm_120a, et à 241 les tests de fumée sur GPU, avec 16 garde-fous automatiques dans le dépôt.

Les prérequis pour cuda-oxide se présentent ainsi :

Composant	cuda-oxide 0.3 (SIMT)	cutile-rs (Tile)
Rust	1.99.1 stable, fixé, avec rustc-dev et llvm-tools	Stable 1.89 ou supérieur
CUDA	Toolkit 13.0 ou supérieur et pilote 13.x (R580+)	13.3 recommandé
Compilation	En amont, Rust vers PTX	JIT, Rust vers Tile IR et cubin
Autres	cargo-oxide, Clang 21 pour bindgen ; LLVM 21 optionnel	GPU sm_80 ou supérieur

La documentation avertit que le llc de LLVM 20 et des versions antérieures ne gère pas les intrinsèques de TMA, tcgen05 et WGMMA émis par le compilateur. Et NVIDIA elle-même tempère les attentes : le README du dépôt le définit comme un projet de recherche à un stade précoce, avec des bugs, des fonctionnalités incomplètes et des changements d’API prévisibles, et sa page pour développeurs le qualifie d’expérimental.

Aucune donnée ne montre non plus que le Rust compilé avec cuda-oxide surpasse CUDA C++. Les notes évoquent des améliorations de code qui visent, selon notre traduction, à « réduire l’écart avec du CUDA C++ réglé à la main », ce qui laisse entendre que cet écart existe. Le projet cohabite par ailleurs avec des initiatives communautaires comme rust-cuda, qui compile Rust vers NVVM IR, et Rust-GPU, orienté SPIR-V et graphisme ; le livre de cuda-oxide affirme que NVIDIA collabore avec les mainteneurs de rust-cuda et présente les deux projets comme complémentaires.

Questions fréquentes

CUDA Oxide 0.3 fonctionne-t-il déjà avec Rust stable ?

Oui. Il utilise Rust stable 1.99.1, mais fixe cette version et nécessite les composants rustc-dev et llvm-tools, car le backend s’appuie sur des parties internes de rustc.

Quelle différence entre cuda-oxide et cutile-rs ?

cuda-oxide suit le modèle SIMT et compile vers PTX avec un contrôle direct des threads et de la mémoire partagée. cutile-rs utilise le modèle CUDA Tile, compile à l’exécution et laisse au compilateur la répartition du travail. Depuis la 0.3, ils partagent un runtime d’hôte.

cute-rs fonctionne-t-il sur les GPU Hopper comme le H100 ?

Non. Dans la version 0.3.0, il ne prend en charge que sm_100a (classe B200) et sm_120a (RTX 50), et rejette Hopper avant la compilation.

Est-il prêt pour la production ?

NVIDIA le décrit comme expérimental et à un stade précoce, et met en garde contre des bugs et des changements d’API. Les démonstrations prouvent qu’il fonctionne, pas qu’il est plus performant que CUDA C++.

Sources :