Aller au contenu

6 · Mojo

Le pari de Modular : un seul langage, de la boucle de service HTTP jusqu'à l'instruction tensor core. Syntaxe Python, performance C++, et un compilateur bâti sur MLIR.


6.1 La proposition

Aujourd'hui, une pile d'inférence typique empile :

Python          ← orchestration, API
C++             ← moteur d'inférence
CUDA C++        ← noyaux
PTX / assembleur← ajustements

Quatre langages, quatre chaînes d'outils, quatre modèles mentaux, et des frontières où l'information de performance se perd.

Mojo propose un seul langage couvrant toute la pile, avec une syntaxe familière aux utilisateurs de Python et un système de types et de métaprogrammation suffisant pour écrire des noyaux.


6.2 À quoi ça ressemble

from gpu import thread_idx, block_idx, block_dim, barrier
from gpu.host import DeviceContext
from memory import UnsafePointer

fn vecteur_add(
    a: UnsafePointer[Float32],
    b: UnsafePointer[Float32],
    c: UnsafePointer[Float32],
    n: Int,
):
    var i = block_idx.x * block_dim.x + thread_idx.x
    if i < n:
        c[i] = a[i] + b[i]

fn main() raises:
    with DeviceContext() as ctx:
        var a = ctx.enqueue_create_buffer[DType.float32](1024)
        # ...
        ctx.enqueue_function[vecteur_add](
            a.unsafe_ptr(), b.unsafe_ptr(), c.unsafe_ptr(), 1024,
            grid_dim=4, block_dim=256,
        )
        ctx.synchronize()

Le modèle est reconnaissable : thread_idx, block_idx, barrier(). Mojo n'invente pas un nouveau modèle de programmation GPU ; il en fournit une implémentation dans un langage unifié.

Au-dessus, il propose des abstractions de plus haut niveau : LayoutTensor et TileTensor, qui jouent un rôle analogue aux layouts de CuTe.


6.3 Les arguments techniques

Argument Détail
Un seul langage pas de FFI, pas de pybind11, pas de duplication de logique
MLIR de bout en bout le compilateur voit tout, du modèle au noyau
Métaprogrammation à la compilation @parameter, spécialisation par cible sans templates C++
Portabilité NVIDIA, AMD, et Apple Silicon (via Metal) pour une partie
Interopérabilité Python appel de bibliothèques Python existantes

Le troisième point mérite d'être développé. Là où CUDA C++ exige des templates pour spécialiser par architecture, Mojo permet :

@parameter
if has_nvidia_gpu_accelerator():
    # chemin wgmma
elif has_amd_gpu_accelerator():
    # chemin MFMA

évalué à la compilation, sans surcoût.


6.4 Les structured kernels

Modular a formalisé une organisation en trois composants pour ses noyaux :

┌──────────────────────────────────────────┐
│ Couche de configuration partagée         │  tailles de tuiles, précisions
├──────────────────────────────────────────┤
│ Couche de gestion des données partagée   │  layouts, mouvements, TMA
├──────────────────────────────────────────┤
│ Trois composants :                        │
│   · producteur (chargement)               │
│   · calcul (MMA)                          │
│   · épilogue                              │
└──────────────────────────────────────────┘

C'est exactement la structure décrite en CUDA moderne · Spécialisation des warps, rendue explicite par le langage. Modular annonce « performance de pointe pour la moitié du code ».


6.5 Les GPU Puzzles

La ressource pédagogique la plus intéressante de l'écosystème Mojo, et recommandable même si vous n'utilisez pas Mojo :

Mojo GPU Puzzles — 34 exercices progressifs, chacun un petit noyau à compléter, validé par des tests automatiques.

La progression : manipulation mémoire bas niveau d'abord, puis transition vers les abstractions TileTensor. Les puzzles couvrent les motifs réels de Triton, FlashAttention et des noyaux de production.

Point pratique notable : la plupart des puzzles fonctionnent sur Apple Silicon via Metal, ce qui les rend accessibles sans carte NVIDIA. C'est rare et précieux pour apprendre.

C'est l'équivalent moderne des GPU Puzzles de Sasha Rush (en Python/NumPy), avec l'avantage de s'exécuter réellement sur un GPU.


6.6 L'évaluation honnête

Ce qui plaide pour Mojo

  • L'unification de la pile est un vrai problème, et personne d'autre ne l'attaque de front.
  • La portabilité NVIDIA + AMD + Apple est plus large que celle de CUDA.
  • Le matériel pédagogique (GPU Puzzles) est excellent.
  • MLIR est le bon choix d'infrastructure, celui de Triton et de CUDA Tile.

Ce qui plaide contre

  • L'écosystème. Trente ans de bibliothèques scientifiques sont en C, C++ et Python. Mojo interopère avec Python, mais l'inertie est massive.
  • Le langage évolue. La syntaxe et les bibliothèques changent encore ; un code écrit aujourd'hui demandera des ajustements.
  • Un fournisseur unique. Modular est une entreprise privée. Le langage a été progressivement ouvert, mais le risque de dépendance reste.
  • Peu de code de production public. L'essentiel de ce qui existe est produit par Modular.

Verdict pratique : intéressant à connaître, excellent pour apprendre la programmation GPU (les puzzles), risqué comme fondation d'un projet à long terme en 2026. À réévaluer.


6.7 Ce que Mojo révèle du domaine

Au-delà du langage lui-même, l'existence de Mojo, de Triton, de Helion, de cuTile et du CuTe DSL raconte la même histoire :

Le CUDA C++ n'est plus considéré comme un point d'arrivée acceptable pour écrire des noyaux.

Cinq projets indépendants, portés par OpenAI, Meta, NVIDIA, Stanford et Modular, convergent vers les mêmes idées : programmer par tuiles, laisser le compilateur gérer les layouts, et rendre les noyaux portables entre générations.

C'est probablement le changement structurel le plus important de la programmation GPU depuis l'introduction des tensor cores.


Résumé du chapitre

À retenir

  • Mojo vise un seul langage de l'orchestration au noyau, avec une syntaxe Python et un compilateur MLIR.
  • Le modèle GPU est celui de CUDA (thread_idx, barrier), avec des abstractions de tuiles (LayoutTensor, TileTensor) au-dessus.
  • Les structured kernels explicitent la décomposition producteur / calcul / épilogue.
  • Les GPU Puzzles sont une excellente ressource d'apprentissage, et fonctionnent en partie sur Apple Silicon.
  • Risques : écosystème jeune, langage en évolution, fournisseur unique.
  • L'existence de Mojo, Triton, Helion, cuTile et CuTe DSL confirme que le modèle par tuiles remplace le modèle par threads.

Vérifiez que vous avez compris

Pourquoi « un seul langage » est-il un argument technique et pas seulement de confort ?

Parce que chaque frontière de langage est une frontière d'optimisation.

Quand du Python appelle du C++ qui appelle du CUDA, aucun compilateur ne voit l'ensemble. Impossible de propager une constante depuis la configuration Python jusqu'à la spécialisation d'un noyau, ou de fusionner deux opérations définies de part et d'autre de la frontière.

Avec une pile unifiée en MLIR, ces optimisations inter-couches deviennent possibles. C'est le même argument qui justifie XLA, torch.compile et les compilateurs de megakernels.

Les GPU Puzzles fonctionnent sur Apple Silicon. Comment, alors que Metal n'a ni warps de 32 ni mémoire partagée au sens CUDA ?

Metal a bien les deux concepts, sous d'autres noms : les SIMD-groups (32 voies sur Apple Silicon) et la threadgroup memory. Le modèle d'exécution est structurellement identique à CUDA.

Ce qui diffère : pas de tensor cores exposés au même niveau (Apple a des unités matricielles mais l'accès direct est limité), une mémoire unifiée avec le CPU, et des tailles de threadgroup différentes.

Les puzzles portent sur les motifs fondamentaux — indexation, mémoire partagée, réductions, tuiles — qui sont communs à tous les GPU. C'est précisément pourquoi ils sont portables.

Vous devez choisir un langage pour un moteur d'inférence à écrire en 2026, à maintenir cinq ans. Mojo ?

Probablement pas comme fondation, pour trois raisons non techniques :

  1. Recrutement — trouver des ingénieurs qui connaissent Mojo est aujourd'hui difficile.
  2. Bibliothèques — vous devrez interfacer NCCL, des formats de modèles, des serveurs HTTP, des systèmes de télémétrie. L'écosystème C++/Python les a tous.
  3. Risque de fournisseur — cinq ans, c'est long pour un langage porté par une seule entreprise.

Le choix conservateur reste C++ pour le moteur, CUDA/Triton pour les noyaux, Python pour l'API. Ce qui n'empêche pas de suivre Mojo : si le pari réussit, la migration sera plus facile depuis une architecture propre que depuis un monolithe.


Chapitre suivant : 7 · Les bibliothèques NVIDIA


Sources de ce chapitre