2 · Triton et Gluon¶

Le point d'entrée réaliste de la programmation GPU en 2026. Ce chapitre couvre Triton en profondeur, puis Gluon — la couche basse qu'OpenAI a dû ajouter quand l'abstraction de Triton a atteint sa limite sur Hopper.

2.1 Le modèle de programmation¶

Triton renverse l'abstraction de CUDA.

	CUDA	Triton
Unité de programme	un thread	un bloc (appelé program)
Ce que vous manipulez	des scalaires	des tuiles (tableaux N-D)
`threadIdx`	visible	inexistant
Mémoire partagée	explicite	gérée par le compilateur
Coalescence	à votre charge	déduite par le compilateur
Tensor cores	dispositions manuelles	émis automatiquement par `tl.dot`

Vous écrivez : « ce programme charge une tuile 128×64, la multiplie, la range ». Le compilateur décide combien de threads, comment les répartir, où mettre la mémoire partagée, et quelles instructions émettre.

2.2 Un premier noyau¶

import torch
import triton
import triton.language as tl

@triton.jit
def add_kernel(x_ptr, y_ptr, sortie_ptr, n, BLOC: tl.constexpr):
    pid = tl.program_id(axis=0)              # identifiant du bloc
    debut = pid * BLOC
    offsets = debut + tl.arange(0, BLOC)     # une tuile d'indices
    masque = offsets < n

    x = tl.load(x_ptr + offsets, mask=masque)
    y = tl.load(y_ptr + offsets, mask=masque)
    tl.store(sortie_ptr + offsets, x + y, mask=masque)


def add(x: torch.Tensor, y: torch.Tensor):
    sortie = torch.empty_like(x)
    n = x.numel()
    grille = lambda meta: (triton.cdiv(n, meta["BLOC"]),)
    add_kernel[grille](x, y, sortie, n, BLOC=1024)
    return sortie

Points à relever :

tl.arange(0, BLOC) crée une tuile, pas une boucle ;
mask= gère les bords sans if ;
BLOC: tl.constexpr est une constante de compilation : chaque valeur produit un noyau compilé distinct ;
la grille est une fonction de meta, ce qui permet à l'autotuner de varier BLOC.

2.3 Le réglage automatique¶

@triton.autotune(
    configs=[
        triton.Config({"BLOC_M": 128, "BLOC_N": 128, "BLOC_K": 32},
                      num_warps=8, num_stages=4),
        triton.Config({"BLOC_M": 128, "BLOC_N": 64,  "BLOC_K": 64},
                      num_warps=4, num_stages=5),
        triton.Config({"BLOC_M": 64,  "BLOC_N": 128, "BLOC_K": 64},
                      num_warps=4, num_stages=4),
    ],
    key=["M", "N", "K"],       # recompile si ces valeurs changent
)
@triton.jit
def matmul_kernel(...): ...

Trois paramètres à comprendre :

Paramètre	Effet
`BLOC_*`	taille des tuiles → intensité arithmétique et mémoire partagée
`num_warps`	nombre de warps par bloc → occupancy
`num_stages`	profondeur du pipeline logiciel → recouvrement contre mémoire partagée

num_stages est le paramètre que Triton expose et que CUDA cache : c'est exactement le nombre d'étages du pipeline cp.async de la partie 4.

L'autotuning coûte cher au premier appel

Avec 20 configurations et 5 tailles distinctes, c'est 100 compilations et 100 séries de mesures. Cela peut prendre plusieurs minutes.

En production : fixez les configurations gagnantes après mesure, ou activez le cache Triton (TRITON_CACHE_DIR). Ne laissez pas l'autotuner tourner à chaque démarrage d'un service.

2.4 Une GEMM en Triton¶

Le noyau canonique, condensé :

@triton.jit
def matmul_kernel(a_ptr, b_ptr, c_ptr, M, N, K,
                  stride_am, stride_ak, stride_bk, stride_bn,
                  stride_cm, stride_cn,
                  BLOC_M: tl.constexpr, BLOC_N: tl.constexpr,
                  BLOC_K: tl.constexpr, GROUPE_M: tl.constexpr):
    pid = tl.program_id(0)

    # Ré-ordonnancement en groupes, pour la localité L2
    n_pid_m = tl.cdiv(M, BLOC_M)
    n_pid_n = tl.cdiv(N, BLOC_N)
    n_par_groupe = GROUPE_M * n_pid_n
    groupe = pid // n_par_groupe
    debut_m = groupe * GROUPE_M
    taille_groupe = min(n_pid_m - debut_m, GROUPE_M)
    pid_m = debut_m + ((pid % n_par_groupe) % taille_groupe)
    pid_n = (pid % n_par_groupe) // taille_groupe

    offs_m = (pid_m * BLOC_M + tl.arange(0, BLOC_M)) % M
    offs_n = (pid_n * BLOC_N + tl.arange(0, BLOC_N)) % N
    offs_k = tl.arange(0, BLOC_K)

    a_ptrs = a_ptr + offs_m[:, None] * stride_am + offs_k[None, :] * stride_ak
    b_ptrs = b_ptr + offs_k[:, None] * stride_bk + offs_n[None, :] * stride_bn

    acc = tl.zeros((BLOC_M, BLOC_N), dtype=tl.float32)
    for k in range(0, tl.cdiv(K, BLOC_K)):
        a = tl.load(a_ptrs, mask=offs_k[None, :] < K - k * BLOC_K, other=0.0)
        b = tl.load(b_ptrs, mask=offs_k[:, None] < K - k * BLOC_K, other=0.0)
        acc = tl.dot(a, b, acc)              # ← émet wgmma/tcgen05
        a_ptrs += BLOC_K * stride_ak
        b_ptrs += BLOC_K * stride_bk

    c = acc.to(tl.float16)
    offs_cm = pid_m * BLOC_M + tl.arange(0, BLOC_M)
    offs_cn = pid_n * BLOC_N + tl.arange(0, BLOC_N)
    c_ptrs = c_ptr + offs_cm[:, None] * stride_cm + offs_cn[None, :] * stride_cn
    tl.store(c_ptrs, c, mask=(offs_cm[:, None] < M) & (offs_cn[None, :] < N))

Deux choses valent d'être notées.

tl.dot fait tout. Il choisit l'instruction MMA (mma, wgmma, tcgen05 selon la cible), la disposition des registres, et le swizzle de la mémoire partagée. C'est l'abstraction la plus précieuse de Triton.

Le regroupement des blocs (GROUPE_M) n'est pas cosmétique. Il change l'ordre de parcours des tuiles pour que les blocs exécutés simultanément partagent des lignes de \(\mathbf{A}\) ou de \(\mathbf{B}\), ce qui améliore le taux de réussite du L2. Gain typique : 10 à 30 %. C'est le même problème que celui que résolvent les tile schedulers de CUTLASS.

2.5 Ce que Triton fait pour vous, en détail¶

Le compilateur Triton effectue, dans l'ordre :

Inférence de layout : quel thread détient quel élément de tuile ;
Allocation de mémoire partagée : quelles tuiles y résident, avec quel swizzle ;
Coalescence : réorganisation des accès pour minimiser les secteurs ;
Pipelining logiciel : émission de cp.async / TMA sur num_stages étages ;
Sélection d'instructions : tl.dot → mma / wgmma / tcgen05 ;
Fusion élémentaire : les opérations sur tuiles restent en registres ;
Spécialisation des warps (progressive, depuis 2025).

Le tout via MLIR, avec plusieurs dialectes intermédiaires (triton, triton_gpu, llvm).

Pour inspecter le résultat :

export MLIR_ENABLE_DUMP=1        # dumps des passes MLIR
export TRITON_CACHE_DIR=./cache  # PTX et cubin générés y atterrissent

Puis cuobjdump -sass cache/**/*.cubin pour vérifier les instructions émises.

2.6 Les back-ends¶

Back-end	État en 2026
NVIDIA (CUDA)	référence, complet
AMD (ROCm)	fonctionnel, supporté par AMD, utilisé en production
Intel (XPU)	actif
CPU	expérimental

La portabilité de Triton est réelle mais pas gratuite : les configurations d'autotuning ne se transposent pas d'un fournisseur à l'autre, et certaines constructions performantes sur NVIDIA le sont moins sur AMD (le wavefront de 64 change les seuils).

2.7 Gluon¶

Depuis 2025, Triton expose un second niveau : Gluon.

Gluon is Triton's lower-level GPU programming model. It exposes layouts, shared memory, warp specialization, and target-specific features directly so advanced kernels can trade convenience for control.

Les deux partagent :

le même frontend Python et la même infrastructure JIT ;
le même modèle de tuiles SPMD.

Ils diffèrent sur ce que le compilateur décide :

	Triton	Gluon
Layouts de tuiles	inférés	explicites
Mémoire partagée	allouée automatiquement	allouée à la main
Mouvement de données	inféré	explicite
Asynchronisme	géré	à votre charge
Spécialisation des warps	partielle, automatique	explicite

# Gluon — syntaxe indicative, voir la documentation officielle
import triton.experimental.gluon.language as gl

@gluon.jit
def kernel(...):
    layout = gl.BlockedLayout(size_per_thread=[1, 8],
                              threads_per_warp=[8, 4],
                              warps_per_cta=[4, 1],
                              order=[1, 0])
    x = gl.load(ptr + offsets, layout=layout)
    smem = gl.allocate_shared_memory(gl.float16, [128, 64], layout=...)
    ...

Pourquoi Gluon existe

Parce que la promesse de Triton — « le compilateur gère les layouts » — se heurte à une réalité : sur Hopper et Blackwell, atteindre le pic exige de contrôler exactement les layouts, le pipelining et la spécialisation des warps.

Plutôt que de dégrader l'abstraction de Triton en y ajoutant des dizaines de paramètres, OpenAI a créé un second langage partageant la même infrastructure. C'est la même stratégie que NVIDIA avec CUTLASS C++ et le CuTe DSL : deux niveaux, une seule pile.

Il existe aussi TLX (Triton Low-level eXtensions), une approche voisine poussée par la communauté, qui ajoute des primitives bas niveau à Triton lui-même plutôt qu'un langage séparé.

2.8 L'écosystème Triton¶

Ce qui est écrit en Triton et que vous pouvez utiliser directement :

Projet	Contenu
Liger Kernel (LinkedIn)	noyaux fusionnés pour l'entraînement de LLM : RMSNorm, RoPE, SwiGLU, CrossEntropy fusionnée
FlashInfer	noyaux d'attention et de service LLM, utilisés par vLLM et SGLang
Unsloth	noyaux d'affinage économe en mémoire
vLLM	de nombreux noyaux internes
PyTorch Inductor	génère du Triton — le back-end par défaut de `torch.compile`

Ce dernier point est important : quand vous appelez torch.compile, vous exécutez du Triton. Savoir le lire aide à comprendre ce que le compilateur a produit :

import os
os.environ["TORCH_COMPILE_DEBUG"] = "1"
# le code Triton généré est écrit dans /tmp/torchinductor_*/

2.9 Les pièges classiques¶

Six erreurs fréquentes en Triton

Oublier le masque dans tl.load / tl.store → lecture hors borne. Triton ne vous protège pas.
BLOC non puissance de 2 → beaucoup d'opérations l'exigent.
Autotuner en production → plusieurs minutes au démarrage.
Confondre tl.constexpr et argument runtime → recompilation à chaque appel si une valeur variable est marquée constexpr.
Accumuler en float16 → tl.zeros(..., dtype=tl.float32) doit être la règle, même pour des entrées en 16 bits.
Oublier le regroupement de blocs dans une GEMM → 10 à 30 % perdus par défaut de localité L2.

Résumé du chapitre¶

À retenir

Triton programme au niveau du bloc et manipule des tuiles. Pas de threadIdx, pas de mémoire partagée explicite.
tl.dot sélectionne automatiquement mma / wgmma / tcgen05.
num_stages est la profondeur du pipeline logiciel ; num_warps l'occupancy ; BLOC_* la taille des tuiles.
Le regroupement de blocs (GROUPE_M) améliore la localité L2 de 10 à 30 % : ce n'est pas optionnel.
Back-ends NVIDIA, AMD et Intel fonctionnels, mais l'autotuning ne se transpose pas.
Gluon expose les layouts, la mémoire partagée et la spécialisation des warps, pour les noyaux où l'abstraction de Triton plafonne.
PyTorch Inductor génère du Triton : savoir le lire aide à comprendre torch.compile.

Vérifiez que vous avez compris¶

Pourquoi tl.zeros((M, N), dtype=tl.float32) même quand les entrées sont en float16 ?

Parce que l'accumulation en 16 bits perd trop de précision. Pour une GEMM de dimension interne \(K = 4096\), l'erreur relative accumulée en FP16 (\(\varepsilon \approx 5\times10^{-4}\)) est de l'ordre de \(\sqrt{4096} \times 5\times10^{-4} \approx 3\ \%\).

C'est aussi ce que font les tensor cores nativement : entrées 16 bits, accumulateur 32 bits. Voir Fondations · Nombres flottants.

Votre noyau Triton est 2× plus lent que le noyau PyTorch équivalent. Que vérifier ?

Dans l'ordre :

La ligne de base — PyTorch appelle peut-être cuBLAS ou cuDNN, qui sont des années-personnes d'optimisation. Un facteur 2 est normal et parfois le bon résultat.
L'autotuning — avez-vous exploré assez de configurations ? Les BLOC_* par défaut sont rarement optimaux.
La mesure — l'échauffement inclut-il la compilation JIT ?
Le regroupement de blocs si c'est une GEMM.
Le PTX généré (TRITON_CACHE_DIR) — les instructions MMA attendues sont-elles présentes ?

Peut-on écrire un megakernel en Triton ?

Non, pas dans l'état actuel. Un megakernel exige :

un noyau persistant dont le nombre de blocs correspond exactement au nombre de SM ;
de la synchronisation inter-blocs par compteurs en mémoire globale ;
une gestion manuelle de la mémoire partagée comme ressource allouée/libérée dynamiquement ;
un interpréteur dispatchant des instructions hétérogènes.

Triton n'expose aucune de ces quatre choses. Gluon en expose une partie (mémoire partagée, spécialisation des warps) mais pas la synchronisation inter-blocs.

C'est pourquoi tous les megakernels de la partie 8 sont en CUDA C++ (souvent via ThunderKittens), et pourquoi les auteurs de Mirage MPK notent que « PyTorch, Triton et TVM ne supportent pas nativement la génération de megakernels de bout en bout ».

Chapitre suivant : 3 · CUTLASS et CuTe DSL