Aller au contenu

6 · CUDA Tile

En août 2025, NVIDIA a fait quelque chose qu'elle n'avait pas fait depuis 2006 : ajouter un second modèle de programmation à CUDA. Ce chapitre explique ce qu'il est, ce qu'il change, et ce qu'il ne fait pas encore.

6.1 Le constat qui l'a motivé

Le modèle SIMT vous fait écrire du code du point de vue d'un thread :

int i = blockIdx.x * blockDim.x + threadIdx.x;
c[i] = a[i] + b[i];

Cela fonctionne remarquablement bien pour des opérations élémentaires. Cela fonctionne mal dès qu'on veut utiliser les tensor cores, parce que ceux-ci opèrent sur des tuiles, avec des dispositions de registres imposées.

Écrire une GEMM performante en SIMT exige donc de :

  • décomposer manuellement le problème en tuiles ;
  • affecter manuellement des fragments de tuile à chaque thread ;
  • respecter des dispositions de registres définies par le PTX ISA ;
  • gérer le swizzling, les descripteurs, les barrières asynchrones.

Autrement dit : le programmeur fait un travail que le compilateur pourrait faire, et le refait à chaque génération de matériel.

Triton avait démontré depuis 2021 qu'un modèle de tuiles est viable. CUDA Tile est la réponse de NVIDIA.

6.2 Le modèle

Dans le modèle Tile, on décrit des opérations sur des tuiles, et le compilateur décide de la répartition sur les threads.

# cuTile Python — syntaxe indicative
import cuda.tile as ct

@ct.jit
def gemm(A, B, C):
    i, j = ct.program_id()
    acc = ct.zeros((128, 128), dtype=ct.float32)
    for k in range(0, A.shape[1], 64):
        a = ct.load(A, (i * 128, k), (128, 64))
        b = ct.load(B, (k, j * 128), (64, 128))
        acc += ct.dot(a, b)
    ct.store(C, (i * 128, j * 128), acc)

Aucun threadIdx. Aucune mémoire partagée explicite. Aucun descripteur. Le compilateur choisit :

  • la taille de bloc ;
  • l'allocation de mémoire partagée et le nombre d'étages de pipeline ;
  • l'émission des copies TMA ;
  • la disposition des données pour wgmma ou tcgen05 ;
  • le swizzling.

Ce n'est pas Triton, mais c'est le même paradigme

La différence principale est l'intégration : CUDA Tile passe par une nouvelle représentation intermédiaire, la CUDA Tile IR, qui est une ISA virtuelle au même titre que PTX. Elle est destinée à devenir une cible de compilation pour d'autres langages, pas seulement pour cuTile Python.

NVIDIA annonce également une implémentation en C++ dans une version ultérieure de CUDA.

6.3 Les trois composants

Composant Statut (août 2026) Rôle
CUDA Tile IR livrée ISA virtuelle par tuiles, sous PTX
cuTile Python livré (CUDA 13.1) DSL Python pour écrire des noyaux par tuiles
Tile en C++ annoncé même modèle, dans CUDA C++

Extensions livrées dans les versions suivantes (13.2 à 13.4) : formats numériques et précision mixte, MMA sur tuiles, vues et sous-vues de tuiles, réductions (y compris atomiques et définies par l'utilisateur), synchronisation et communication au niveau des tuiles.

Le profilage est supporté : Nsight Compute sait profiler les noyaux CUDA Tile depuis CUDA 13.1.

6.4 Le support matériel

Architecture Compute capability Supporté
Ampere 8.x oui
Ada 11.x* oui
Hopper 9.x (non listé dans l'annonce initiale)
Blackwell 10.x, 12.x oui

* La numérotation des compute capabilities a été réorganisée dans CUDA 13 ; se référer à la documentation pour la correspondance exacte de votre carte.

L'absence de Hopper dans la liste initiale est notable

L'annonce de CUDA 13.1 mentionne Ampere, Ada et Blackwell. Hopper — l'architecture la plus déployée en centre de données en 2026 — n'y figure pas explicitement, sans que la raison en soit publique.

Vérifiez le support de votre cible avant de bâtir dessus. La couverture s'élargit à chaque version.

6.5 Ce que ça change, et ce que ça ne change pas

Ce que ça change

1. La portabilité entre générations. Un noyau écrit en tuiles peut être recompilé pour Blackwell puis Rubin sans réécriture. C'est l'argument le plus fort : le code SIMT optimisé pour sm_90a (TMA + wgmma) ne tourne pas sur Blackwell.

2. Le coût d'entrée. Écrire une GEMM correcte en cuTile prend des heures ; en CUDA C++ avec TMA et wgmma, des semaines.

3. L'accès aux fonctionnalités récentes. Le compilateur émet TMA, tcgen05, le swizzling et le pipelining sans que le programmeur ait à les connaître.

Ce que ça ne change pas

1. Le plafond de performance reste celui du compilateur. Comme pour Triton, il y aura toujours des noyaux où l'écriture manuelle gagne. La question est de savoir de combien, et pour quelle proportion de cas.

2. Les problèmes structurels restent. CUDA Tile ne supprime pas les frontières de noyau, ne fait pas de synchronisation inter-blocs, et ne construit pas de megakernel. Il optimise l'intérieur d'un noyau.

3. Il faut toujours comprendre le roofline. Un noyau limité par la mémoire le reste, quel que soit le modèle de programmation.

6.6 Où se situe CUDA Tile dans le paysage

Abstraction ↑
            │  torch.compile / Inductor
            │  Helion
            │  Triton  ·  cuTile Python
            │  Gluon   ·  CuTe DSL      ← contrôle explicite en Python
            │  CUTLASS C++  ·  ThunderKittens
            │  CUDA C++ (TMA, wgmma manuels)
            │  PTX
            ↓  SASS

cuTile Python se place au niveau de Triton, avec deux différences :

  • il est officiel, donc suivi par NVIDIA génération après génération ;
  • il n'est pas portable vers AMD ou Intel, alors que Triton l'est.

C'est le compromis habituel : intégration verticale contre portabilité.

La lecture stratégique

L'existence de CUDA Tile est un aveu implicite : le modèle SIMT ne suffit plus pour exploiter le matériel moderne, et Triton avait raison.

Pour un développeur, la conséquence pratique est rassurante : le modèle de tuiles que vous apprenez avec Triton est transférable. tl.load, tl.dot, tl.store et leurs équivalents cuTile décrivent la même chose.

6.7 Les autres nouveautés de CUDA 13.x qui comptent

Puisque ce chapitre parle du CUDA le plus récent, voici ce qu'il faut savoir d'autre.

Nouveauté Version Utilité
Green Contexts (API runtime) 13.1 partitionner le GPU au niveau des SM
cuBLAS : +2 à 6× sur BF16/FP8/block-scaled Blackwell 13.1 gratuit, recompilez
Grouped GEMM dans cuBLAS 13.1 essentiel pour les MoE
CCCL 3.1 : réductions déterministes, CUB en une phase 13.1 reproductibilité
Compute Sanitizer : détection d'erreurs mémoire à la compilation 13.1 débogage
Abandon de la compilation hors ligne pré-Turing 13.0 vérifiez vos cibles
Fin des pilotes Windows fournis avec le toolkit 13.0 installation séparée

Le point grouped GEMM mérite d'être souligné : c'est exactement ce dont un MoE a besoin (une GEMM par expert, avec des tailles différentes), et son absence obligeait jusque-là à passer par des bibliothèques tierces.

Résumé du chapitre

À retenir

  • CUDA Tile (CUDA 13.1, décembre 2025) est un second modèle de programmation CUDA, à base de tuiles plutôt que de threads.
  • Trois composants : la CUDA Tile IR (ISA virtuelle), cuTile Python (livré), et une implémentation C++ annoncée.
  • Le compilateur choisit la répartition sur les threads, la mémoire partagée, le pipelining, TMA et les instructions MMA.
  • Support : Ampere, Ada, Blackwell. Vérifiez pour Hopper.
  • Il optimise l'intérieur d'un noyau : il ne supprime ni les frontières de noyau ni le besoin de comprendre le roofline.
  • Son existence confirme que le modèle par tuiles, popularisé par Triton, est devenu le modèle de référence.

Vérifiez que vous avez compris

Pourquoi NVIDIA introduit-elle un second modèle plutôt que d'améliorer le compilateur SIMT ?

Parce que l'information manque au compilateur. Dans du code SIMT, la notion de « tuile » n'existe pas : le compilateur voit des indices calculés à partir de threadIdx et des accès mémoire individuels. Reconstruire une structure de tuile à partir de cela est un problème d'analyse très difficile, et fragile.

Dans le modèle Tile, la structure est déclarée. Le compilateur sait qu'il manipule un bloc 128×64 et peut donc choisir librement la disposition, le swizzle, l'instruction MMA. C'est le même argument que celui de Triton.

cuTile ou Triton pour un nouveau projet en 2026 ?

Cela dépend d'un seul critère : allez-vous cibler autre chose que NVIDIA ?

  • Non → cuTile a l'avantage du support officiel, de l'intégration avec Nsight, et du suivi des futures architectures.
  • Oui (AMD, Intel, ou incertitude) → Triton, qui a des back-ends AMD et Intel fonctionnels et un écosystème mature (Liger Kernel, FlashInfer, vLLM, une communauté).

À la date de rédaction, Triton reste le choix par défaut pour la plupart des projets, simplement parce qu'il est éprouvé depuis 2021 et que l'écosystème de noyaux qui l'utilisent est immense. cuTile est jeune.

CUDA Tile peut-il générer un megakernel ?

Non, pas dans son état actuel. CUDA Tile optimise la structure interne d'un noyau : découpage en tuiles, pipelining, choix d'instructions.

Un megakernel exige autre chose : une synchronisation inter-blocs à grain fin, un ordonnancement dynamique de tâches hétérogènes, et une gestion de la mémoire partagée comme ressource globale. Ce sont des problèmes de structure entre les opérations, pas dans une opération.

C'est précisément le vide que comblent Mirage MPK, Event Tensor et les autres compilateurs de la partie 8 — dont les auteurs notent explicitement que « PyTorch, Triton et TVM ne supportent pas nativement la génération de megakernels de bout en bout ».

Partie suivante : Écosystème

Sources de ce chapitre