Caractéristiques du GPU H100
- Offre 80 Go de mémoire HBM3 et une bande passante de 3 To/s (en pratique, elle est un peu plus faible)
- Offre 50 Mo de cache L2 et une bande passante de 12 To/s. Il est divisé en deux sections de 25 Mo sur le GPU et relié par un crossbar (le crossbar est une source de dégradation des performances)
- Composé de 132 Streaming Multiprocessors (SM), chaque SM ayant la configuration suivante :
- Jusqu’à 227 Ko de mémoire partagée dans un cache L1 de 256 Ko (soit une bande passante combinée d’environ 33 To/s)
- Fournit un Tensor Memory Accelerator (TMA) permettant la génération d’adresses asynchrone et le fetching mémoire. Il offre aussi des fonctions comme la prise en charge du réseau mémoire on-chip, mais elles ne sont pas abordées dans ce billet
- Divisé en 4 quadrants, chacun composé d’un warp scheduler, de 512 vector registers (chacun contenant 32 mots de 4 octets), d’un tensor core pour la multiplication de matrices, et d’instructions intégrées prenant en charge des opérations parallèles comme sum/multiply
Conseils d’optimisation des performances du GPU H100
- Il est essentiel d’exploiter au maximum les Tensor Cores. Le H100 offre 989 TFLOPs de performances en multiplication de matrices FP16, donc le taux d’utilisation global du GPU dépend fortement du niveau d’exploitation des Tensor Cores
- Toutefois, pour exploiter au maximum les Tensor Cores, il faut prendre en compte les points suivants :
- L’utilisation de l’instruction Warp Group Matrix Multiply Accumulate (WGMMA) est indispensable, mais sa mise en œuvre est délicate
- La Shared Memory n’est pas aussi rapide qu’on pourrait le penser, et son utilisation demande beaucoup de soin
- La génération d’adresses a un coût élevé, il faut donc optimiser avec des mécanismes comme le Tensor Memory Accelerator (TMA)
- Augmenter l’occupancy reste utile, et les registres constituent la ressource principale
- Ces caractéristiques s’appliquent dans une certaine mesure à d’autres GPU que le H100, mais l’exploitation des Tensor Cores est particulièrement importante et délicate sur le H100
ThunderKittens : une DSL embarquée CUDA optimisée pour le H100
- Une DSL embarquée basée sur CUDA, développée pour tirer le maximum des performances des GPU récents comme le NVIDIA H100
- Fournit les 4 types de templates basés sur des tuiles ci-dessous :
- Register Tile (tenseur 2D stocké dans les registres)
- Register Vector (tenseur 1D stocké dans les registres)
- Shared Tile (tenseur 2D stocké dans la Shared Memory)
- Shared Vector (tenseur 1D stocké dans la Shared Memory)
- Propose aussi divers opérateurs de manipulation de tuiles (
exp, mul, sum, etc.) au niveau warp ou groupe de warps
- La réécriture des kernels existants de Flash Attention et Flash Attention 2 avec ThunderKittens simplifie fortement le code et apporte jusqu’à 30 % de gain de performances sur le H100
- Un kernel Based Linear Attention a également été implémenté avec ThunderKittens, atteignant 215 TFLOPs (et plus de 300 TFLOPs en incluant le recompute, du fait des caractéristiques de l’algorithme)
Réflexion d’un point de vue philosophique
- Si ThunderKittens fonctionne bien, c’est parce qu’il ne cherche pas à tout prendre en charge et propose à la place une abstraction simple, basée sur des tuiles, qui correspond bien à l’architecture des GPU
- Utiliser des registres vectoriels de 1024 bits au lieu des mots traditionnels de 32 bits est une avancée, mais cela demande surtout un changement de paradigme : considérer la tuile 16x16 comme unité de registre
- Étant donné que les workloads IA tournent au final surtout autour de la multiplication de matrices, des reductions et du reshape, l’approche basée sur des tuiles est pertinente ; du point de vue matériel, on peut aussi s’attendre à une évolution non seulement autour des systolic arrays, mais aussi vers la prise en charge de petites multiplications de matrices
- Il faut aller plus loin et changer de façon de penser pour concevoir les algorithmes d’IA dans une forme optimisée pour le matériel. Par exemple, limiter la taille d’état d’un RNN pour qu’elle tienne dans un SM, et ajuster la densité de calcul au niveau exigé par le matériel
L’avis de GN⁺
- ThunderKittens peut être une option séduisante pour les développeurs déjà familiers avec CUDA, car il permet d’obtenir facilement des gains de performances sans modifier en profondeur le code des kernels existants
- Pour les débutants, la barrière à l’entrée peut toutefois rester élevée. Il faudra sans doute davantage d’exemples de code et de ressources d’apprentissage à l’avenir
- Le projet d’étendre la prise en charge de ThunderKittens à d’autres GPU, notamment ceux d’AMD, est intéressant. Cela pourrait contribuer à réduire la dépendance à un fournisseur
- En fin de compte, concevoir les modèles et algorithmes d’IA eux-mêmes sous une forme optimisée pour le matériel est un point extrêmement important. Cela suppose d’abord une compréhension approfondie des caractéristiques du matériel, et ThunderKittens peut aider les développeurs à acquérir ce type d’insight
- On peut s’attendre à ce que les efforts continus de R&D et les contributions open source de Hazy Research dynamisent fortement l’écosystème CUDA. À plus long terme, l’apparition de frameworks à un niveau d’abstraction plus élevé semble toutefois aussi nécessaire
1 commentaires
Avis Hacker News
Les exigences du matériel IA deviennent de plus en plus claires. Les GPU ont été conçus à l’origine pour d’autres usages, mais ils sont utilisés pour l’IA parce qu’ils disposent d’un bon matériel pour la multiplication matricielle. Les « GPU IA » pourraient supprimer certaines fonctionnalités des GPU classiques, et la tendance va vers des nombres plus courts (virgule flottante 16 bits, 8 bits, 2 bits et 1 bit). Cet article suggère que le matériel qui aime les tuiles 16x16 présente de nombreux avantages.
Il faut des coprocesseurs dédiés à l’IA (NPU), en particulier dans les systèmes desktop prosumer destinés aux développeurs, experts et joueurs. Les GPU fonctionnent en entreprise, mais ils sont peu pratiques pour l’IA du point de vue de l’informatique personnelle. Les limites de la VRAM et surtout l’absence d’API ouvertes standard en dehors de Vulkan posent problème.
Pour faire progresser la recherche en IA, il faut mieux étudier les neurosciences et la psychologie. Des éléments liés à la topologie des graphes des réseaux neuronaux pourraient également être pertinents.
Est-ce que c’est CUTLASS rendu plus convivial ?
La mascotte de ThunderKittens a une vibe chat / Sony Aibo. On dirait qu’elle a été bien générée par une IA.
Si l’Universal Basic Compute (UBC) devait être envisagé comme substitut au revenu universel de base, ce serait un futur très dystopique. Imaginez une seule entreprise comme Nvidia produire toute la puissance de calcul.