CUDA Tile : programmer le GPU à la tuile, plus au thread

Pourquoi écrire un kernel GPU est resté un métier d’expert

Le matériel qui fait le calcul d’un LLM a changé de nature, mais le modèle pour le programmer, lui, a peu bougé. L’essentiel du travail d’inférence est de la matmul matmul Multiplication de matrices, l'opération dominante des couches d'un réseau de neurones. Quand on dit qu'un GPU « calcule », il fait à 90 % des matmul. exécutée par les Tensor Cores, des unités qui opèrent sur des blocs de matrice, pas sur des scalaires isolés. Pourtant, écrire un kernel CUDA CUDA Compute Unified Device Architecture. La plateforme de calcul GPU de NVIDIA : langage, compilateur et bibliothèques (cuBLAS, cuDNN). Son écosystème logiciel est le principal verrou face aux alternatives comme ROCm ; à l'exécution, son « contexte » réserve aussi une part incompressible de VRAM. Approfondir dans le glossaire revient encore, depuis quinze ans, à raisonner à l’échelle du thread : découper les données, assigner à chaque thread sa trajectoire, orchestrer à la main la mémoire partagée et les synchronisations.

Cet écart (un matériel qui pense en tuiles, un langage qui fait penser en threads) est exactement ce que CUDA Tile entend combler. Et il faut comprendre d’abord le modèle qu’il remonte d’un cran, sinon la nouveauté reste abstraite.

Le modèle SIMT : ce que CUDA Tile abstrait

CUDA, dans sa forme historique, est un modèle SPMD SPMD Single Program, Multiple Data. Modèle de parallélisme où tous les rangs exécutent exactement le même programme sur des fragments de données différents. C'est le modèle pour lequel le TPU est taillé ; sa version divergente, le MPMD (un rang qui fait du logging en plus, par exemple), cassait l'optimisation. à grain thread, baptisé SIMT (Single Instruction, Multiple Threads). Vous écrivez le code d’un thread ; le GPU en lance des dizaines de milliers, regroupés par warps de 32 qui exécutent la même instruction en parallèle sur des données différentes. C’est puissant, et c’est laborieux. Pour qu’un kernel de matmul soit rapide, le programmeur doit : découper les matrices en sous-blocs qui tiennent dans la mémoire partagée d’un SM SM Streaming Multiprocessor. Bloc de calcul indépendant d'un GPU NVIDIA, contenant ses propres unités d'exécution, registres, mémoire partagée et Tensor Cores. Un H100 SXM5 en compte 132. L'occupation des SM est la métrique-clé de saturation d'un GPU. Approfondir dans le glossaire , charger ces blocs en coordonnant les threads pour éviter les conflits de banc mémoire, synchroniser les warps aux bons moments (__syncthreads()), puis, sur les générations récentes, émettre explicitement les instructions Tensor Core (MMA, matrix multiply-accumulate) et gérer les accélérateurs de transfert mémoire. Chacune de ces étapes est une source de bug et un levier de performance, et toutes dépendent de l’architecture cible.

Résultat : le kernel performant est l’apanage d’une poignée d’experts, et il est souvent à réécrire à chaque génération de GPU. C’est un coût, de temps et de talents rares, qui pèse sur tout l’écosystème, NVIDIA compris.

La tuile comme unité de programmation

CUDA Tile déplace le curseur. La description qu’en donne NVIDIA mérite d’être citée mot pour mot, parce qu’elle dit tout le changement :

Concrètement, vous ne pilotez plus les threads : vous chargez une tuile, vous calculez dessus comme sur un tableau, vous la rangez. La parallélisation devient l’affaire du compilateur (« the compiler and runtime determine the best way to launch that work onto individual threads »). Et l’usage du matériel spécialisé n’est plus à votre charge (« cuTile will leverage the advanced capabilities of NVIDIA hardware (such as tensor cores, shared memory, and tensor memory accelerators) without requiring explicit programming »).

cuTile, le langage dédié en Python, donne à voir le contraste. Un additionneur de vecteurs s’écrit sans une seule mention de thread :

# cuTile : API en preview, extrait illustratif
import cuda.tile as ct

@ct.kernel
def vector_add(a, b, c, tile_size: ct.Constant[int]):
    pid     = ct.bid(0)                                   # identifiant de bloc
    a_tile  = ct.load(a, index=(pid,), shape=(tile_size,))
    b_tile  = ct.load(b, index=(pid,), shape=(tile_size,))
    result  = a_tile + b_tile                             # calcul sur la tuile entière
    ct.store(c, index=(pid,), tile=result)

Pas de threadIdx.x, pas de __shared__, pas de __syncthreads(). Là où la version SIMT exigerait de calculer un indice global par thread et de borner les débordements, la version tuile manipule a_tile comme un objet entier. « You can focus on your algorithm at a higher level, while the NVIDIA CUDA compiler and runtime handle partitioning your tile algorithm into threads and launching them onto the GPU », résume NVIDIA. Le DSL est d’abord sorti en Python (CUDA 13.1), puis en C++ avec CUDA 13.3 (mai 2026).

Ce n’est pas une idée neuve en soi : c’est précisément le modèle de Triton, le DSL tuile d’OpenAI devenu la lingua franca des kernels custom de l’écosystème PyTorch. La nouveauté, c’est que NVIDIA l’adopte dans CUDA, et le fait reposer sur une brique de plus bas niveau qui, elle, est la vraie pièce d’architecture.

Tile IR : un ISA virtuel, et l’astuce de la compatibilité

Sous cuTile, il y a Tile IR, et NVIDIA le qualifie sans détour : « a new virtual instruction set architecture (ISA) for programming NVIDIA GPUs ». Le terme ISA virtuel est la clé de tout, et il a un précédent que tout développeur CUDA connaît : le PTX.

Le PTX est l’ISA virtuel à grain thread de CUDA. Quand vous compilez un kernel, vous ne produisez pas directement le code machine d’un GPU précis (le SASS) : vous produisez du PTX, et le pilote l’abaisse (le JIT) vers le SASS de la carte présente au moment de l’exécution. C’est ce niveau d’indirection qui permet à un binaire compilé en 2020 de tourner sur un GPU sorti en 2025. Tile IR applique exactement cette idée, un étage plus haut : vous exprimez votre calcul en tuiles, Tile IR le capture, et le compilateur l’abaisse vers l’architecture cible. La promesse qui en découle est explicite dans la documentation : « your tile code will be compatible with future GPU architectures ».

Pourquoi un ISA virtuel à grain tuile, et pourquoi maintenant ? Parce que l’unité de calcul utile a migré. Tant que le GPU était un océan de threads scalaires, le PTX était le bon niveau d’abstraction portable. Depuis que les Tensor Cores (qui consomment des tuiles) font l’essentiel du travail, capturer l’intention au niveau de la tuile laisse au compilateur la liberté de choisir l’instruction MMA, la taille de bloc et le schéma mémoire optimaux pour chaque architecture, là où un kernel SIMT fige ces choix dans le code source. Tom’s Hardware le formule comme une réorganisation du modèle logiciel CUDA autour de l’exécution tensor-native, la base posée pour Blackwell, puis Rubin et Feynman.

Triton, i4 et l’effet écosystème

Le pari de NVIDIA ne s’arrête pas à fournir un DSL de plus. Deux signaux montrent qu’il s’agit de capter tout l’écosystème de l’écriture de kernels, pas seulement les nouveaux venus.

Le premier : NVIDIA a publié un backend Tile IR pour OpenAI Triton. Triton est le DSL tuile dominant, celui que le compilateur inductor de PyTorch émet pour les kernels générés. Lui donner un chemin d’abaissement vers Tile IR, c’est faire passer les kernels Triton existants par la même infrastructure que cuTile. NVIDIA joue donc sur deux tableaux : cuTile comme frontend maison, et Tile IR comme socle commun sous le frontend que la communauté utilise déjà. Quel que soit le langage par lequel vous entrez, vous ressortez par l’IR de NVIDIA.

Le second : le support du type i4 (entier 4 bits), ajouté à Tile IR non pas au lancement mais dans une mise à jour ultérieure, CUDA 13.3 (mai 2026), pour la quantification quantification Réduction du nombre de bits codant chaque poids d'un modèle (de 16 bits vers 8, 4, voire moins). Elle divise l'empreinte mémoire d'autant, au prix d'une perte de précision contrôlée, sans changer le nombre de paramètres. Approfondir dans le glossaire . Les tuiles i4 doivent être converties vers un type entier supporté avant calcul, et les opérations de conversion et d’empaquetage (exti, trunci, pack, unpack) ont été étendues pour le gérer. C’est l’aveu, au niveau de l’IR, que la basse précision n’est plus un cas marginal mais un format de premier ordre : le pendant entier des formats flottants FP4 FP4 Format à virgule flottante 4 bits, frontière 2026 de l'inférence à haut débit. Quatre fois moins de mémoire que le FP16, mais une portée dynamique très étroite : ne tient qu'avec un scaling fin via formats à blocs (MXFP4, NVFP4). Approfondir dans le glossaire dont chaque runtime choisit déjà le dialecte.

Ce que ça change, et pour qui

Vous écrivez des kernels custom pour de l’inférence ou de l’entraînement. C’est le cœur de cible. CUDA Tile promet de réduire le temps passé à orchestrer des threads pour le réinvestir dans l’algorithme, et de ne pas tout réécrire à la prochaine architecture. Le réflexe sain reste le banc d’essai : NVIDIA n’a pas publié de comparatif de débit face à un kernel CUDA C++ taillé à la main, et c’est précisément le chiffre qui décide d’une migration. Mesurez sur votre opérateur avant de remplacer un kernel critique. Un signal encourageant côté recherche : des travaux indépendants commencent à publier des implémentations cuTile d’opérateurs exigeants comme FlashAttention FlashAttention Implémentation tile-by-tile de l'attention qui évite de matérialiser la matrice d'attention complète en HBM. Réduit drastiquement la consommation mémoire et accélère le calcul, particulièrement sur longs contextes. Trois versions (v1/v2/v3), chacune optimisée pour une génération de GPU. Approfondir dans le glossaire sur matériel Blackwell, preuve que le modèle tient sur des charges réelles, pas seulement sur l’additionneur de vecteurs.

Vous êtes déjà sur Triton. Le backend Tile IR vous concerne sans changer votre code : c’est un nouveau chemin d’abaissement, à évaluer pour ce qu’il apporte en performance et en couverture d’opérateurs sur les GPU récents.

Vous opérez un parc hétérogène ou plus ancien. Tempérez l’enthousiasme. Le compilateur vise d’abord Blackwell ; sur Ampere ou Ada, le support arrive par étapes, et la promesse de portabilité ascendante reste à éprouver. Pour un kernel qui doit tourner aujourd’hui sur tout votre parc, le couple CUDA C++ / PTX reste la valeur sûre.

Une dernière nuance de cadrage. Le billet d’annonce de CUDA 13.1 met en avant des gains de 2× à 6×, mais ce sont des accélérations de cuBLAS sur Blackwell face au H200, pas des mesures de CUDA Tile. Ne reportez pas ces chiffres sur le modèle de programmation : ils décrivent les bibliothèques de la version, pas le langage à tuiles.

Conclusion

L’abstraction de la programmation GPU monte d’un cran, et ce n’est pas un hasard de calendrier. Le PTX a tenu une décennie parce que le thread était la bonne unité portable d’un GPU généraliste. Tile IR fait le pari que la tuile est la bonne unité d’un GPU devenu une machine à Tensor Cores, et que le même code source pourra viser Blackwell aujourd’hui, Rubin et Feynman demain, sans la réécriture qui rythmait jusqu’ici chaque transition.

C’est aussi un geste sur le terrain où NVIDIA défend sa position la plus précieuse, qui n’est pas le silicium mais le logiciel : le sujet de fond de la bataille du compilateur que Google attaque côté TPU, et de l’écart CUDA / ROCm. En adoptant le modèle tuile et en se branchant sous Triton, NVIDIA réduit la friction d’écriture de kernels à l’intérieur de son écosystème au moment précis où les alternatives essaient d’en abaisser le coût de sortie. La question ouverte n’est plus « peut-on programmer le GPU autrement qu’au thread » : la réponse est livrée. C’est : à quelle distance du kernel écrit à la main le compilateur tuile saura-t-il rester, une fois passé l’additionneur de vecteurs et confronté aux opérateurs qui, eux, paient vraiment la facture mémoire ?

Sources et méthode

Cet article s’appuie sur un travail de vérification arrêté au 16 juin 2026. Étiquettes : fait vérifié = source primaire citable ; estimation crédible = cohérent mais non vérifié indépendamment ; hypothèse = raisonnement sans mesure.

CUDA Tile, cuTile et Tile IR

• NVIDIA Technical Blog, NVIDIA CUDA 13.1 Powers Next-Gen GPU Programming with NVIDIA CUDA Tile : developer.nvidia.com (fait vérifié). Citations verbatim sur le modèle SIMT vs tuile, « virtual instruction set architecture (ISA) », « compatible with future GPU architectures », et l’abstraction des Tensor Cores. Les gains 2×–6× y concernent cuBLAS sur Blackwell vs H200, pas CUDA Tile (clarification éditoriale).
• NVIDIA Technical Blog, Simplify GPU Programming with NVIDIA CUDA Tile in Python : developer.nvidia.com (fait vérifié). Exemple vector_add (ct.kernel, ct.load, ct.store), prérequis (compute capability 8.x/10.x/11.x/12.x, pilote R580+/R590, CUDA 13.1+, Python 3.10+).
• NVIDIA Technical Blog, Develop High-Performance GPU Kernels in C++ with NVIDIA CUDA Tile (26 mai 2026) : developer.nvidia.com, implémentation C++ livrée avec CUDA 13.3 (fait vérifié).
• Page produit CUDA Tile, paquet cuda-tile (PyPI) et dépôt NVIDIA/cutile-python (fait vérifié, existence et nature).
• Limites du compilateur tileiras (Blackwell d’abord) en CUDA 13.1, et type i4 ajouté en CUDA 13.3 (27 mai 2026) : release notes CUDA 13.x (fait vérifié). Annonce de CUDA 13.1 début décembre 2025 ; couverture presse : Phoronix.

Backend Triton et contexte

• NVIDIA Technical Blog, Advancing GPU Programming with the CUDA Tile IR Backend for OpenAI Triton : developer.nvidia.com (fait vérifié quant à l’existence du backend).
• Cadrage tensor-native et perspective Rubin/Feynman : Tom’s Hardware (source secondaire, analyse).
• Implémentation cuTile de FlashAttention sur GB10 : arXiv 2601.16032 (publication de recherche tierce ; signal d’usage réel, gains non repris ici).

Points à manier avec prudence

• Aucune mesure tête-à-tête publiée de CUDA Tile vs kernel CUDA C++ écrit à la main, ni vs Triton, au 16 juin 2026. L’argumentaire NVIDIA porte sur la productivité et la portabilité, pas sur un gain de débit chiffré.
• Compatibilité ascendante : présentée comme objectif (« compatible with future GPU architectures »), pas comme garantie démontrée ; le compilateur cible d’abord Blackwell.
• API en preview. L’extrait cuTile (import cuda.tile as ct, noms de fonctions) est illustratif : l’API évolue, vérifiez la syntaxe exacte sur la documentation de votre version.

Image d’en-tête

• Image d’en-tête. Wafer de silicium photographié par Rob Bulmahn, Wikimedia Commons, sous licence CC BY 2.0, recadré pour l’en-tête, la grille régulière des puces du wafer servant de métaphore du pavage (tiles) au cœur de CUDA Tile.