Aller au contenu

1 · Panorama des langages

Le tableau de décision, et le raisonnement qui le fonde.

1.1 Le gradient complet

Niveau Outil Ce que vous écrivez Ce que le compilateur fait
8 torch.compile du PyTorch tout : fusion, choix de noyaux, CUDA Graphs
7 Helion des boucles sur des tuiles découpage, réglage automatique, génère du Triton
6 Triton des opérations sur des tuiles disposition, mémoire partagée, pipelining, MMA
5 Gluon / TLX des tuiles + layouts explicites émission des instructions
4 CuTe DSL des layouts et des atomes ordonnancement fin, en Python
3 CUTLASS C++ des composants assemblés rien, mais tout est fourni
3 ThunderKittens des tuiles 16×16 enveloppe TMA/wgmma
2 CUDA C++ des threads allocation de registres
1 PTX des instructions virtuelles mapping vers SASS
0 SASS rien (lecture seule) —

1.2 Les quatre critères de décision

Critère 1 — Quelle est la nature de votre opération ?

Nature Outil naturel
Élémentaire, fusion de plusieurs ops Triton, ou torch.compile
Réduction, normalisation, softmax Triton
Produit matriciel générique cuBLAS (ne rien écrire)
Produit matriciel avec épilogue ou format inhabituel CUTLASS / CuTe DSL
Attention (variante standard) FlashAttention (ne rien écrire)
Attention (variante custom) Triton, ThunderKittens
Convolution standard cuDNN
Structure de données irrégulière (graphe, tri, hachage) CUDA C++
Enchaînement de dizaines d'opérations megakernel → CUDA C++

La première question est toujours : existe-t-il une bibliothèque ?

La majorité des noyaux écrits à la main dans l'industrie n'auraient pas dû l'être. cuBLAS, cuDNN, CUB, FlashAttention, FlashInfer, CUTLASS couvrent l'essentiel des opérations standard, mieux que ce que vous écrirez.

Écrivez un noyau quand votre opération est inhabituelle — une fusion spécifique, une forme non standard, un format de quantification maison.

Critère 2 — Quelles cibles matérielles ?

Cible Options
NVIDIA seul tout
NVIDIA + AMD Triton, SYCL, Kokkos, Mojo, HIP (via HIPIFY)
+ Intel Triton, SYCL, Mojo
+ Apple Mojo (partiel), MLX, Metal direct
+ navigateur WebGPU
Toutes aucune, réellement

Critère 3 — Quel est votre budget de temps ?

Ordres de grandeur pour une GEMM performante, à compétence égale :

Outil Temps à ~80 % du pic
cuBLAS 5 minutes
Triton 1 à 3 jours
CuTe DSL 1 à 2 semaines
CUTLASS C++ 2 à 4 semaines
CUDA C++ brut avec TMA/wgmma 1 à 3 mois

Ces chiffres sont indicatifs et supposent une familiarité préalable avec l'outil. Le premier projet dans chaque outil coûte deux à trois fois plus.

Critère 4 — Qui maintiendra ce code ?

Un noyau CUDA C++ avec TMA, wgmma et des barrières asynchrones est illisible pour quiconque n'a pas ce contexte. Un noyau Triton se lit comme du NumPy.

C'est un critère sérieux, pas une considération de confort : la moitié des noyaux optimisés à la main deviennent non maintenables dès que leur auteur change d'équipe.

1.3 Comparaison sur un exemple

Un softmax sur les lignes d'une matrice, dans quatre outils.

PyTorch

y = torch.softmax(x, dim=-1)

Appelle un noyau cuDNN/ATen. Optimal pour les formes courantes.

Triton

import triton
import triton.language as tl

@triton.jit
def softmax_kernel(sortie_ptr, entree_ptr, stride, n_cols, BLOC: tl.constexpr):
    ligne = tl.program_id(0)
    cols  = tl.arange(0, BLOC)
    masque = cols < n_cols

    x = tl.load(entree_ptr + ligne * stride + cols,
                mask=masque, other=-float('inf'))
    x = x - tl.max(x, axis=0)
    num = tl.exp(x)
    y   = num / tl.sum(num, axis=0)
    tl.store(sortie_ptr + ligne * stride + cols, y, mask=masque)

Une vingtaine de lignes, pas de gestion de threads, pas de mémoire partagée explicite, et les réductions tl.max / tl.sum sont générées efficacement.

CUDA C++

Environ 60 lignes, avec les réductions de warp explicites, la mémoire partagée pour les partiels, le déroulage — la version développée dans la partie 2.

CUDA C++ optimisé pour Hopper

150 à 300 lignes : chargement TMA, pipeline à plusieurs étages, spécialisation des warps, réductions par cluster. Nettement plus rapide sur les grandes lignes, et écrit par très peu de gens.

Le rapport effort/gain, concrètement

Sur un softmax de dimension raisonnable, Triton atteint typiquement 80 à 95 % de ce que fait une implémentation CUDA experte, pour 10 % du temps de développement.

L'écart se creuse sur les opérations où la disposition des données est critique — les GEMM et l'attention. C'est exactement pourquoi Gluon et le CuTe DSL existent : donner le contrôle des layouts sans quitter Python.

1.4 Ce que Triton ne sait pas faire

Il faut être précis sur les limites, parce qu'elles déterminent quand descendre.

Limitation Conséquence Alternative
Pas de contrôle des layouts de registres plafond sur les GEMM de pointe Gluon, CuTe DSL
Spécialisation des warps limitée pipelining moins agressif Gluon, CUTLASS
Pas de synchronisation inter-blocs pas de megakernel CUDA C++
Pas de contrôle du swizzle conflits résiduels CuTe
Compilation JIT au premier appel latence de démarrage mise en cache
Autotuning parfois long itération lente fixer les configs

Le point « pas de synchronisation inter-blocs » est le plus structurant pour ce document : aucun megakernel n'est écrit en Triton, et c'est pour cela que la partie 8 revient au CUDA C++.

1.5 Recommandation par profil

Si vous faites du deep learning appliqué

  1. torch.compile en premier réflexe ;
  2. Triton pour vos fusions spécifiques ;
  3. les bibliothèques (FlashAttention, FlashInfer, cuBLAS) pour tout le reste ;
  4. CUDA C++ seulement si vous construisez un système d'inférence.

Si vous faites de la recherche en systèmes ML

  1. CUDA C++ et ThunderKittens — vous aurez besoin du contrôle ;
  2. CuTe DSL pour les GEMM et l'attention ;
  3. Triton pour prototyper vite ;
  4. lisez CUTLASS comme référence de conception.

Si vous faites du calcul scientifique

  1. CUDA C++ si NVIDIA seul, avec Thrust/CUB ;
  2. SYCL ou Kokkos si le parc est mixte ;
  3. OpenMP target ou OpenACC pour porter du Fortran existant ;
  4. attention au FP64 : vérifiez le débit de votre cible.

Si vous faites du graphique ou de l'embarqué

  1. Vulkan compute si vous êtes déjà dans Vulkan ;
  2. WebGPU pour le navigateur ;
  3. Metal sur Apple ;
  4. CUDA sur Jetson.

1.6 Le tableau de synthèse

Outil Cibles Niveau Maturité (2026) Point fort
CUDA C++ NVIDIA bas très mûre contrôle total
Triton NVIDIA, AMD, Intel moyen mûre rapport effort/perf
Gluon NVIDIA, AMD moyen-bas jeune layouts explicites en Python
CuTe DSL NVIDIA bas bêta GEMM de pointe en Python
CUTLASS C++ NVIDIA bas très mûre référence de la GEMM
ThunderKittens NVIDIA bas active lisibilité des tuiles
Helion NVIDIA, AMD haut bêta autotuning, portabilité
Mojo NVIDIA, AMD, Apple moyen en évolution langage unique
cuTile NVIDIA moyen jeune officiel, portable entre générations
HIP AMD (+ NVIDIA) bas mûre portage CUDA
SYCL tous moyen mûre standard, C++ moderne
OpenCL tous bas héritage portée matérielle maximale
Vulkan compute tous bas mûre intégration graphique
WebGPU navigateurs moyen mûre en 2026 déploiement web
Metal Apple bas mûre seule option Apple native

Résumé du chapitre

À retenir

  • Le gradient va de torch.compile à SASS ; descendez d'un cran seulement quand le profileur le justifie.
  • La première question est toujours : existe-t-il une bibliothèque ? Vous ne battrez pas cuBLAS, cuDNN ou FlashAttention.
  • Triton est le point d'entrée réaliste : ~90 % de la performance experte pour ~10 % de l'effort, sur les fusions et les réductions.
  • Triton ne fait pas de synchronisation inter-blocs : pas de megakernel en Triton.
  • Le critère de maintenabilité est réel : un noyau wgmma manuel devient non maintenable dès que son auteur part.

Vérifiez que vous avez compris

Vous devez écrire une attention avec un motif de masque exotique. Par quoi commencer ?

Par vérifier que FlexAttention (PyTorch) ou FlashInfer ne le couvrent pas : les deux prennent des fonctions de masque arbitraires et génèrent le noyau.

Si vraiment non : Triton. L'attention est le cas d'usage pour lequel il a le meilleur rapport effort/performance, et le tutoriel officiel de FlashAttention en Triton est un excellent point de départ.

ThunderKittens ou CuTe DSL seulement si la version Triton se révèle insuffisante après mesure — et il faudra alors mesurer contre cuDNN, pas contre votre première version.

Votre laboratoire a un parc mixte H100 et MI300X. Un seul code source. Quelles options ?

Trois, avec des compromis différents.

  1. Triton : back-ends NVIDIA et ROCm fonctionnels. Le plus simple si votre code est du deep learning. Attention : les configurations d'autotuning ne se transposent pas, il faut régler par cible.
  2. HIP : écrit une fois, compile pour les deux (hipcc cible aussi NVIDIA). Bas niveau, et le portage des optimisations Hopper (TMA, wgmma) n'a pas d'équivalent AMD.
  3. SYCL : le plus standard, mais l'écart de performance mesuré au niveau du noyau est de 10 à 30 % contre du code natif.

En pratique, la plupart des équipes maintiennent deux chemins optimisés et un chemin portable de repli.

Pourquoi OpenAI a-t-elle créé Gluon en plus de Triton, plutôt que d'améliorer Triton ?

Parce que les deux objectifs sont contradictoires. Triton promet que le compilateur gère les layouts, la mémoire partagée, le mouvement de données et l'asynchronisme. C'est son intérêt et sa limite.

Sur Hopper et Blackwell, atteindre le pic exige de contrôler exactement ces choses — spécialisation des warps, layouts de tensor cores, allocation de mémoire partagée. Les exposer dans Triton reviendrait à détruire son abstraction.

Gluon partage le frontend, la JIT et le modèle de tuiles SPMD de Triton, mais expose « layouts, mémoire partagée, spécialisation des warps et fonctionnalités spécifiques à la cible ». C'est le même arbitrage que CUTLASS C++ contre CuTe DSL : deux niveaux, une seule pile.

Chapitre suivant : 2 · Triton et Gluon

Sources de ce chapitre