Comment faire tourner un GPU AMD à la vitesse du « brrr »

• HipKittens est un ensemble de primitives de programmation conçu pour exploiter les performances potentielles des GPU AMD, en optimisant l’accès mémoire, l’ordonnancement et la réutilisation du cache
• Le GPU AMD MI355X adopte une architecture avec 256 compute units et 8 chiplets (XCD), et offre un grand fichier de registres ainsi que des instructions de cœur matriciel fines
• Contrairement à NVIDIA, AMD ne dispose pas de réattribution de registres, d’instructions matricielles asynchrones ni de mbarrier, ce qui rend l’ordonnancement 8-wave ping-pong et 4-wave interleave plus efficace que la wave specialization
• HipKittens améliore la localité des caches L2 et LLC grâce à un ordonnancement de grille conscient des chiplets, et obtient des gains de bande passante et de TFLOPS sur les calculs GEMM et Attention
• Cette approche compense le manque de maturité logicielle de l’écosystème GPU AMD et fournit une base pour accroître l’évolutivité du calcul IA sur des matériels variés

---

Architecture et caractéristiques de performance des GPU AMD CDNA

• Le GPU AMD MI355X comprend 256 compute units (CU), chaque CU étant composée de 4 SIMD
  • Un SIMD exécute une wave composée de 64 threads, à comparer au warp NVIDIA de 32 threads
• Le MI355X dispose d’environ 70 % de la SRAM d’un B200 (165 KB) et ne prend pas en charge les fonctions instructions de multiplication matricielle asynchrones, réattribution de registres, accélération de la mémoire tensorielle, mbarrier
• En revanche, il offre un fichier de registres deux fois plus grand et 60 % de processeurs en plus (256 CU contre 160 SM)
  • Il prend en charge des instructions de cœur matriciel petites et fines, ainsi qu’un chargement direct mémoire globale → mémoire partagée (similaire à TMA)
• AMD adopte une architecture en chiplets composée de 8 chiplets (XCD), chaque XCD disposant de son propre cache L2, avec au-dessus un cache LLC
• D’après le tableau, le MI355X atteint 2,5 PFLOPs en BF16, 5,0 PFLOPs en MXFP8, 10,1 PFLOPs en MXFP6, avec 288 Go de mémoire et 8 To/s de bande passante

Défis de conception de kernels pour AMD

• Optimisation des accès mémoire : à cause des contraintes du compilateur HIPCC et d’un comportement d’E/S non public, la conception de la disposition des données et des motifs de swizzle est essentielle
• Ordonnancement à l’intérieur du processeur : sur AMD, il faut exploiter le fichier de registres et les petites instructions matricielles plutôt que la mémoire partagée
• Ordonnancement entre processeurs : du fait de la structure en chiplets, il faut répartir les tâches en tenant compte des effets NUMA au niveau du cache

Motifs d’accès mémoire de HipKittens

• HipKittens (HK) utilise le tile comme unité de base des données et fournit des fonctions d’opération proches de PyTorch
  • Un tile est défini par le type de données, la taille et le layout, avec prise en charge de différentes entrées via la métaprogrammation par templates C++
• Ordonnancement des registres : comme HIPCC ne peut pas utiliser certains registres comme entrées MFMA, HK fournit une fonction explicite de fixation des registres
  • Le développeur peut attribuer lui-même les registres pour écrire des kernels atteignant les performances maximales
• Disposition des registres : sur AMD, le layout varie selon le type de données et la forme de la matrice, ce qui empêche l’usage d’un motif de swizzle unique
  • Par exemple, un tile bf16 16×16 et un tile bf16 16×32 nécessitent des motifs de swizzle différents
• Structure en phases des instructions : les instructions de mémoire partagée d’AMD possèdent des groupes de phases discontinus et sont peu documentées en interne
  • HK fournit pour cela un solveur obtenu par rétro-ingénierie
• Génération d’adresses : AMD prend en charge le chargement asynchrone HBM → mémoire partagée et effectue des optimisations via le swizzle d’adresses HBM

Ordonnancement intra-processeur : motifs Wave

• La wave specialization est efficace chez NVIDIA, mais sur AMD elle dégrade les performances à cause de l’absence de réattribution des registres
  • Les waves productrices occupent des registres inutiles, tandis que les waves consommatrices subissent des spills faute de registres suffisants
• D’après les expériences de HK, la wave specialization entraîne sur AMD une baisse de l’intensité arithmétique et un goulot d’étranglement mémoire
  • Exemple : sur GEMM, la configuration HK 0/8 atteint 1605 TFLOPs, contre 1570 TFLOPs pour CUTLASS
• Motifs d’ordonnancement alternatifs
  • 8-wave ping-pong : deux waves exécutent en alternance des clusters mémoire/calcul
  • 4-wave interleave : une wave alterne finement mémoire et calcul
  • Le 8-wave produit un code plus concis, tandis que le 4-wave est plus fin mais plus long
  • Sur GEMM et Attention Forward, le 8-wave atteint des performances au niveau de l’état de l’art

Ordonnancement inter-processeur : approche consciente des chiplets

• L’AMD MI355X possède 8 chiplets XCD, chacun avec un cache L2 indépendant
  • Les blocs de threads sont attribués aux chiplets en round-robin, de sorte que l’ordre de la grille influence directement l’efficacité de réutilisation du cache
• Une disposition simple en row-major réduit le taux de réutilisation du cache L2 et provoque une perte de bande passante
  • Exemple : L2 55 %, LLC 95 %, 15,1 To/s, 1113 TFLOPs
• HK introduit un ordonnancement de grille conscient des chiplets afin d’exploiter simultanément la localité des caches L2 et LLC
  • Les blocs de threads sont regroupés par zones adjacentes de la matrice de sortie afin de maximiser la réutilisation des données d’entrée

Exemples de kernels réels

• Les hot loops des kernels Attention Forward et BF16 GEMM utilisent le planning 8-wave ping-pong de HK
  • Chaque boucle exécute alternativement des clusters calcul–mémoire et se synchronise au moyen de barrières d’ordonnancement
  • Dans les exemples de code, des opérations HK comme mma_AtB, load, exp2, col_sum sont utilisées de manière répétée

Conclusion : AMD à l’ère de l’IA multi-silicium

• HipKittens atteint des performances compétitives sur AMD CDNA3 et CDNA4
  • Trois éléments clés : accès mémoire optimisé, ordonnancement wave centré sur AMD et ordonnancement de grille conscient des chiplets
• Les kernels HK atteignent les meilleures performances côté AMD et se montrent compétitifs face aux kernels NVIDIA Blackwell
• Pour la diversité du calcul IA, il est nécessaire d’élargir l’accessibilité aux GPU AMD, et HipKittens fournit pour cela une base logicielle essentielle
• L’amélioration de l’ordonnancement des registres dans HIPCC est désignée comme un axe de progrès important pour la suite

[GN⁺]

Commentaires sur Hacker News

• Je recommande de consulter la discussion connexe sur HipKittens
• Il y a aussi un article sur la même recherche, HipKittens: Fast and furious AMD kernels, avec des commentaires de George Hotz et d’employés d’AMD
• C’est bien de voir le monde académique s’attaquer à ce problème, mais au final je pense que c’est à AMD de le résoudre en interne
  • À mon avis, les entreprises de matériel devraient se concentrer sur le matériel. Ainsi, les incitations restent pures. Même avec 20 % de performances en moins, je pense que ce serait préférable
  • Tout à fait d’accord. AMD a repoussé ce problème il y a 10 ans et essaie seulement maintenant de rattraper son retard. Le matériel est excellent, mais il n’exploite pas son potentiel à cause d’un manque de compétences en écriture de firmware
  • Mais cette équipe de recherche a déjà créé un logiciel similaire pour les GPU Nvidia. On dirait simplement que d’excellents chercheurs mettent leur expertise à profit
  • D’après ce que je sais, AMD traite déjà ce problème à plusieurs niveaux et collabore aussi avec tinycorp
• En lisant l’article, on a l’impression que l’optimisation est difficile à cause de la complexité architecturale des GPU AMD. Mais à long terme, l’approche d’AMD pourrait mieux passer à l’échelle. Là où Nvidia utilise 2 chiplets, AMD en a 8, ce qui crée des problèmes de localité mémoire. Comme le nombre de chiplets augmentera sans doute encore à l’avenir, l’expérience acquise aujourd’hui dans la gestion de cette complexité pourrait s’avérer utile à long terme
  • AMD n’a pas besoin de warp specialization pour obtenir de hautes performances, ce qui simplifie la programmation
• Beaucoup de développeurs ont essayé de faire « go brrr » avec les GPU AMD pour le grand public, mais ont échoué. Je ne comprends pas pourquoi AMD ne résout pas elle-même les problèmes logiciels. Ils ont largement les moyens aujourd’hui, donc ne pas embaucher des développeurs n’est plus une excuse. Leurs GPU pour datacenter ne sont pas mauvais, mais pour les expérimentations personnelles en ML et en IA, Nvidia reste de loin meilleur. J’ai l’impression que ma RTX 3090 vieille de 5 ans reste meilleure que n’importe quel GPU grand public AMD sorti jusqu’ici
  • L’expérience développeur chez AMD est horrible. Ils n’acceptent même pas les rapports de bugs de crash de pilote
  • J’ai récemment remplacé une NVidia 5090 par deux AMD R9700 32GB pour un serveur d’inférence, et l’expérience a été entièrement positive. Ça a fonctionné immédiatement sur le kernel Fedora sans configuration DKMS, et connecter des conteneurs avec ROCm a été facile. Il suffisait de modifier la configuration d’Ollama et de Storyteller. C’était une expérience bien plus agréable que CUDA
  • Nvidia maintient même directement un fork d’Unreal Engine. AMD est à un niveau qui ne permet même pas de rivaliser
  • Nvidia est la seule entreprise de matériel à offrir aux ingénieurs logiciel une rémunération compétitive. Chez AMD, il reste une culture où le logiciel n’est toujours pas vu comme le « vrai travail », et ce genre d’inertie est difficile à changer
• Mojo avait des idées pour améliorer l’expérience développeur (devX) sur les GPU AMD ; je me demande où cela en est
• Je ne comprends pas qu’AMD n’investisse pas des milliards de dollars pour améliorer son logiciel. Nvidia est l’entreprise la plus valorisée du monde, et AMD est son seul concurrent
  • AMD fait des efforts, mais je pense qu’il est difficile de transformer une culture d’entreprise organisée autour d’un renouvellement annuel du matériel en une culture centrée sur le logiciel. Le logiciel ne génère pas de revenus immédiats comme le matériel, donc la direction a tendance à lui accorder une priorité moindre. En plus, le fait que des fournisseurs externes livrent du code en open source peut sembler positif à court terme, mais nuit à la qualité sur le long terme. Il suffit aussi de rater une seule tendance matérielle pour risquer d’être distancé par la concurrence
  • J’ai travaillé chez plusieurs fournisseurs de GPU, et Nvidia est la seule entreprise qui considère le logiciel comme un actif (asset) et y investit. Les autres ne le voient que comme un coût
• Personnellement, je n’aime pas trop le mème « go brr », mais c’est amusant de le voir utilisé dans un endroit comme Stanford
  • En fait, ils utilisaient déjà « go brr » il y a un an lors de la présentation de ThunderKittens
  • Si ce genre de mème apparaît sur les canaux officiels d’une université, c’est peut-être déjà le signe que la mode est passée
• Le projet lui-même est excellent, mais on peut se demander pourquoi AMD ne fait pas cela directement. On a l’impression qu’AMD ne comprend toujours pas l’importance d’une pile logicielle mature. Il lui faut une pile unifiée utilisable sur toutes les cartes, comme CUDA. J’ai longtemps cru qu’AMD finirait par rattraper son retard, mais j’ai presque abandonné cette idée désormais
• Le projet est bon, mais l’article lui-même donne une impression étrangement mal rédigée
  • Le texte est trop maladroit. On dirait qu’il dépend trop de l’IA, ou qu’il imite un style d’écriture généré par IA. Des formulations comme « consultez la partie one » ou « comment faire go brr avec les GPU AMD » reviennent sans cesse. J’ai aussi trouvé dommage que certains points techniques qui auraient dû être expliqués avec des graphiques soient déroulés en 100 lignes de code