- HipKittens est un projet qui fournit des kernels haute performance pour GPU AMD et des primitives de programmation en C++ afin d’améliorer l’efficacité des calculs d’IA
- Dans l’écosystème AMD existant, AITER, PyTorch, Triton, TileLang, Composable Kernel et d’autres ont montré leurs limites en raison de performances instables et d’un support encore immature
- HipKittens s’appuie sur une abstraction fondée sur les tuiles (tile abstraction), en maintenant une interface commune entre NVIDIA et AMD tout en séparant les implémentations propres à chaque matériel
- Des kernels écrits en moins d’environ 500 lignes de code obtiennent de meilleures performances que les kernels AMD existants écrits manuellement en assembleur
- Le projet propose une base concrète pour étendre les calculs d’IA à un environnement multi-silicium et ouvre la possibilité d’une transition vers un écosystème matériel ouvert
État actuel et problèmes de l’écosystème GPU AMD
- Les calculs d’IA ont jusqu’à présent évolué autour d’un seul fournisseur matériel, mais les GPU AMD offrent, sur les générations les plus récentes, des performances de calcul et une bande passante mémoire de tout premier plan
- Exemple : AMD MI355X OAM fournit BF16 2.5 PFLOPs, MXFP8 5.0 PFLOPs, 288 Go de mémoire, 8.0 To/s de bande passante
- Toutefois, en l’absence d’une pile logicielle mature, ces performances ne sont pas pleinement exploitées dans les workflows d’IA réels
- Les principaux composants logiciels d’AMD comprennent AITER, PyTorch, Triton, Mojo, TileLang, Composable Kernel (CK), entre autres
- Les kernels de rétropropagation SDPA Llama GQA de AITER et PyTorch n’atteignent respectivement qu’environ 30 % et 24 % des performances SoTA
- Mojo reste limité à environ 50 % des performances à cause de bank conflicts
- TileLang ne prend en charge que jusqu’à CDNA3, et sa dépendance à CK ainsi que l’absence de fonctions clés augmentent la complexité
- Triton rencontre des difficultés sur le suivi de la durée de vie des registres et l’optimisation des accès mémoire
- En conséquence, les meilleurs kernels AMD doivent encore être écrits directement en assembleur par des experts, ce qui impose de fortes limites en matière de passage à l’échelle et de maintenance
Comparaison avec l’écosystème centré sur CUDA
- Dans la communauté IA, on considère qu’il existe un fossé à l’entrée lié à l’écosystème CUDA (CUDA moat)
- Par le passé, le développement de kernels NVIDIA exigeait lui aussi des années d’efforts sur des bases bas niveau comme CUDA/CUTLASS
- Plus récemment, les avancées des DSL (langages spécialisés) et des outils assistés par IA ont simplifié le développement de kernels NVIDIA
- Exemples : ThunderKittens (TK), CuTe DSL, TinyGrad “tinykittens”, TileLang, Gluon
- C’est dans cette dynamique qu’est explorée, côté AMD, la nécessité de nouvelles primitives de programmation
Principes de conception de HipKittens
- HipKittens est la version développée pour AMD, après ThunderKittens (NVIDIA) et ThunderMittens (Apple Silicon)
- Concept central : l’abstraction fondée sur les tuiles (tile abstraction)
- Universalité de l’abstraction par tuiles – le modèle de calcul fondé sur les tuiles, efficace sur NVIDIA, s’applique naturellement aussi à AMD
- Implémentation backend spécialisée selon l’architecture – les schémas d’accès mémoire et l’ordonnancement des registres sont conçus différemment selon le matériel
- Adaptabilité de la stratégie d’ordonnancement – sur les CDNA3·CDNA4 d’AMD, l’ordonnancement fondé sur les waves est inefficace, mais l’ordonnancement par tuiles conserve simplicité et performances
- En séparant l’interface (tuiles et opérations) de l’implémentation (mapping matériel), le projet propose un modèle commun applicable à diverses architectures GPU
Résultats de performance
- Kernel Attention Forward : écrit en environ 500 lignes de code, il obtient de meilleures performances que le kernel assembleur d’AITER
- Il domine dans les configurations Causal/Non-Causal sur différentes dimensions de tête (D) et longueurs de séquence (N)
- Kernel GEMM : construit autour d’une boucle centrale de moins de 100 lignes, il atteint les meilleures performances en BF16 comme en FP8
- Des améliorations ont aussi été observées sur les kernels Attention Backward, Rotary, Fused Dropout-Residual-LayerNorm par rapport aux meilleures performances existantes
- Tous ces kernels conservent lisibilité et facilité de modification tout en atteignant des performances au niveau de l’assembleur écrit à la main
Extension vers une IA multi-silicium
- Pour réaliser pleinement le potentiel de l’IA, il faut pouvoir exploiter un matériel diversifié et ouvert
- HipKittens vise à faire des GPU AMD une plateforme réellement accessible aux développeurs IA
- En combinaison avec la pile logicielle open source d’AMD, le projet ouvre la voie à un écosystème IA fondé sur le multi-silicium
- Un article de suivi abordera la structure technique détaillée de HipKittens (annonce de la partie two)
1 commentaires
Avis Hacker News
Nous avons un contrat avec AMD et nous entraînons Llama 405B sur MI350X pour MLPerf.
La situation d’AMD s’améliore clairement. Si vous avez un GPU AMD, je recommande d’installer PyTorch via pytorch.org en cliquant sur Linux+ROCm. Il y a 3 ans, c’était désespérant, mais aujourd’hui la plupart des fonctionnalités principales marchent bien. J’ai lancé nanochat sur MI300X et ça a fonctionné immédiatement. MI350X est tout aussi stable.
Bien sûr, AMD reste en retrait par rapport à NVIDIA, et il faut encore beaucoup investir dans l’écosystème logiciel, les compilateurs et les pilotes. Mais comparé à la situation désespérante d’il y a 2 ans, il y a maintenant de l’espoir.
Si vous voulez voir ce qui manque au backend LLVM d’AMD, comparez le code de HipKittens à celui de CUDA Kittens.
Pour l’entraînement, NVIDIA et Google sont n°1, AMD est n°2, et derrière il n’y a pratiquement personne. Intel et Tenstorrent sont encore loin, Huawei faisait planter les exemples par segfault, Groq a renoncé à vendre des puces, Cerebras est introuvable, et j’ai perdu tout intérêt pour Trainium quand il a fallu 5 jours pour obtenir une instance.
La configuration initiale reste encore un peu rugueuse, mais c’est bien mieux qu’avant. Par exemple, il y a seulement un mois, nanochat marchait mal, alors que maintenant ça fonctionne bien. Le point important, c’est que les gens s’intéressent enfin à l’écosystème AMD.
L’IA a besoin de diversité matérielle. Qu’une seule entreprise monopolise tout le matériel et tout le logiciel de l’IA, c’est peut-être bon pour les actionnaires, mais néfaste pour le progrès technologique.
Je ne comprends pas la valorisation de NVIDIA. Aujourd’hui, quelques algorithmes comme les Transformers ont gagné, et les données sont devenues plus importantes. Comme on n’a plus besoin d’une grande diversité de codes HPC comme avant, les concurrents n’ont plus qu’à optimiser un ensemble étroit d’algorithmes.
Si on peut faire tourner efficacement vLLM et les Transformers, un énorme marché s’ouvre. Du coup, j’ai l’impression qu’AMD ou Huawei pourraient rattraper NVIDIA assez facilement, alors je me demande où est son véritable moat. InfiniBand suffit-il à lui seul ?
NVIDIA reste très fort sur les GPU de datacenter, mais Google utilise déjà massivement les TPU, et OpenAI commande aussi du matériel AMD.
L’écosystème CUDA reste important, mais tout le monde cherche activement à s’en émanciper. AMD, Intel, Qualcomm et d’autres se sont aussi lancés dans cette compétition. HipKittens ressemble à une étape importante vers un écosystème logiciel neutre.
AMD pourrait très bien financer ce genre de projet, mais à mon avis ils ont encore raté l’occasion. Avec les GPU et l’IA, c’est toujours la même histoire.
HipKittens est une amélioration, mais en interne AMD manque de capacité pour suivre les performances des kernels. Le DevOps a été sous-traité à TCS en Inde, et eux comprennent mal la situation.
Les équipes avec de bons responsables font tourner leur propre shadow IT. ROCm n’avait pas ce type d’équipe, et les améliorations n’ont commencé qu’après les plaintes de gros clients cloud.
Même lorsqu’ils recrutent des talents, ils proposent des salaires inférieurs au marché. Ils se comparent à des entreprises comme Qualcomm ou Walmart.
Depuis 4 ans, les bonus n’ont jamais été versés intégralement.
Comme avec Radeon VII auparavant, AMD a coupé le support ou bouleversé l’écosystème trop souvent, sans le laisser mûrir. À cause du manque de compatibilité CUDA et d’investissement, la plupart des organisations ont simplement choisi NVIDIA.
Cela dit, AMD investit désormais régulièrement dans ROCm et l’écosystème Instinct, donc la situation s’améliore peu à peu. Mais sur le réseau, NVIDIA reste encore très largement devant.
Je me demande s’il existe des projets construits au-dessus de ThunderKittens.
Si c’est une version portée vers HIP, cela semble avoir une valeur pratique bien plus grande qu’un simple portage CUDA.
L’expression « raw assembly vs cooked assembly » est intéressante.
À une époque, écrire directement du code assembleur pour CPU était courant. Il ne faut pas trop en avoir peur sur GPU non plus. Après tout, c’est bien ce code que le compilateur finit par produire.
D’un point de vue mathématique, l’inférence repose sur des opérations de base d’algèbre linéaire / BLAS.
Je me demande s’il serait possible d’avoir une API d’inférence concise couvrant 80 % des cas d’usage, en tenant compte des dtype et de la sparsity. Elle n’aurait sans doute pas besoin d’être aussi complexe que CUDA.
composable-kernel est le logiciel qui a le plus souvent provoqué des OOM (manques de mémoire) sur mon système Gentoo. Quand Clang compile CK, il utilise plus de 2,5 Go par thread.
Sur les serveurs de build officiels, il faudra probablement réduire le nombre d’architectures. D’après ce que j’ai compris, la plupart des paquets dépendants n’ont besoin que des headers. Des travaux sont en cours pour réduire le temps de compilation.
device_grouped_conv2d_fwd_xdl_ngchw_gkcyx_ngkhw_f16_instance, je ne sais vraiment pas ce qu’ils font subir à clang.Je me demande si ces progrès permettront aussi de bien faire tourner des LLM sur les puces AMD grand public.
Par exemple, j’envisage un HP OMEN MAX 16 avec CPU/iGPU AMD et RTX 5080 ; est-ce que la partie AMD peut rivaliser avec la RTX ?
Cet article de blog ou les résultats avancés sur Mac sont intéressants.
C’est un peu plus lent que l’API Sonnet 4, mais pour des performances locales de ce niveau, j’en suis largement satisfait.
La combinaison Opencode + Ollama + Qwen3 Coder est une excellente alternative à ClaudeCode + Sonnet4.
Évidemment, si vous attendez de l’IA qu’elle fasse tout le codage à votre place, vous le vivrez peut-être différemment. Mais comme assistant personnel, c’est parfait.
Je ne comprends pas pourquoi AMD a complètement ignoré le B300.