Hopper (microarchitecture)

NVIDIA Hopper

Marque déposée 🛈

Caractéristiques

Date de sortie	20 septembre 2022
Procédé	TSMC 4N

Interfaces supportées

Historique

Variantes	Ada Lovelace (grand public)
Prédécesseur	Ampere
Successeur	Blackwell

Crédit image:

Lynn Gilbert

licence CC BY-SA 4.0 🛈

Hopper est une microarchitecture de processeur graphique (GPU) développée par Nvidia. Elle est conçue pour les serveurs des centres de données et est sortie parallèlement à la microarchitecture Ada Lovelace.

Nommée d'après l'informaticienne et contre-amiral de l'United States Navy Grace Hopper, l'architecture Hopper a fuité en novembre 2019 et a été officiellement révélée en mars 2022. Elle améliore les performances de ses prédécesseurs, les microarchitectures Turing et Ampere, possédant un nouveau multiprocesseur de flux et un sous-système mémoire plus rapide.

Le GPU Hopper H100 est fabriqué avec le procédé TSMC 4N et comporte 80 milliards de transistors. Il possède jusqu'à 144 multiprocesseurs de flux^[1]. Dans un socket SXM5 , le Hopper H100 offre de meilleures performances que le PCIe^[2].

Les multiprocesseurs de flux (SM) d'Hopper améliorent ceux des microarchitectures Turing et Ampere, bien que nombre maximal de warps concurrents par SM reste identique entre les architectures Ampere et Hopper, 64^[3]. L'architecture Hopper fournit un Tensor Memory Accelerator (TMA), qui supporte le transfert mémoire asynchrone bidirectionnel entre la mémoire partagée et la mémoire globale^[4]. Avec le TMA, les applications peuvent transférer jusqu'à des tenseurs 5D. Lors d'écritures entre la mémoire partagée et la mémoire globale, les opérateurs de réduction par élément et par bit peuvent être utilisés, évitant les instructions sur registres et SM, tout en permettant aux utilisateurs d'écrire des codes spécialisés dans les warps. Le TMA est activé par cuda::memcpy_async^[5].

Lorsqu'ils parallélisent des applications, les développeurs peuvent utiliser des clusters (groupes) de blocs de threads . Les blocs de threads peuvent effectuer des instructions « atomiques » dans la mémoire partagée d'autres blocs de threads au sein de son groupe, appelée également mémoire partagée distribuée . La mémoire partagée distribuée peut être utilisée par un SM simultanément avec le cache L2 ; lorsqu'elle est utilisée pour transmettre des données entre les SM, elle peut utiliser la bande passante combinée de la mémoire partagée distribuée et du cache L2. Le nombre maximal des groupes « portables » dans un cluster est de 8, bien que le Hopper H100 puisse supporter une nombre maximal de groupes de 16 en utilisant la fonction cudaFuncAttributeNonPortableClusterSizeAllowed, potentiellement au prix d'une réduction du nombre de blocs actifs^[6]. Avec le L2 multicasting et la mémoire partagée distribuée, la bande passante requise pour les lectures et écritures en mémoire vive dynamique est réduite^[7].

Hopper améliore le throughput en format flottant simple précision (FP32), avec deux fois plus d'opérations FP32 par cycle et par SM que son prédécesseur. De plus, l'architecture Hopper possède de nouvelles instructions, dont l'algorithme de Smith-Waterman^[6]. Comme sur Ampere, l'arithmétique TensorFloat-32 (TF-32) est supportée. La grille de formats supportés par les deux architectures est identique^[8].

Le Hopper H100 supporte les mémoires HBM3 et HBM2 jusqu'à 80 GB ; le système de mémoire HBM3 supporte 3 TB/s, en augmentation par rapport à Ampere A100. Au sein de l'architecture, la capacité du cache L2 et la bande passante ont été augmentées^[9].

Hopper permet aux noyaux de calcul des CUDA d'utiliser la compression en ligne, y compris en allocation mémoire individuelle, bien que cette fonctionnalité ne réduise pas l'empreinte mémoire. Le compresseur choisira automatiquement entre plusieurs algorithmes de compression^[9].

Un exemple complet d'allocation de mémoire compressible peut être trouvé à the CUDA samples Github's cudaCompressibleMemory example^[10].

Le Hopper H100 accroît la capacité du cache combiné L1, du cache de texture et de la mémoire partagée à 256 kB. Comme ses prédécesseurs, il combine les caches L1 et de texture dans un cache unifié conçu comme un « buffer coalescent ». L'attribut cudaFuncAttributePreferredSharedMemoryCarveout peut être utilisé pour définir le découpage du cache L1. Hopper introduit des améliorations à NVLink par le biais d'une nouvelle génération ayant une largeur de bande de transmission globalement plus rapide^[11].

Les applications CUDA peuvent rencontrer des interférences lorsqu'elles réalisent des opérations fence/flush à cause de l'ordonnancement de la mémoire . Parce que le GPU ne peut pas savoir lors de l'exécution quelles écritures sont garanties et quelles sont celles qui sont visibles seulement par le hasard du timing, il peut attendre pour des opérations mémoire inutiles, ralentissant donc fortement les opérations fence/flush. Par exemple, lorsqu'un noyau réalise des calculs dans la mémoire GPU tandis qu'un noyau parallèle effectue des communications avec un pair, le noyau local nettoiera (flush) ses écritures, se traduisant par des écritures NVLink ou PCIe plus lentes du noyau de communication. Avec l'architecture Hopper, grâce à la mise en oeuvre des domaines de synchronisation mémoire, le GPU peut réduire le le délai inutile par une opération de fence^[12].

L'interface de programmation (API) mathématique de l'architecture Hopper propose des fonctions dans les SM telles que __viaddmin_s16x2_relu, qui exécute la fonction sur demi-mot ${\displaystyle max(min(a+b,c),0)}$. Dans l'algorithme de Smith-Waterman, la fonction __vimax3_s16x2_relu peut être utilisée, un max sur trois valeurs suivi d'un max avec zéro^[13]. De même, Hopper accélère l'implémentation de l'algorithme de Needleman-Wunsch^[14].

L'architecture Hopper utilise un moteur de transformeur^[15].

Le socket SXM5 du H100 a un thermal design power (TDP) de 700 watts. En regard à son asynchronisme, l'architecture Hopper pourrait atteindre un fort taux d'utilisation et donc avoir une meilleure performance par watt^[16].

En novembre 2019, un célèbre compte Twitter a posté un tweet révélant que l'architecture succédant à Ampere serait appelée Hopper, nommée d'après l'informaticienne et contre-amiral United States Navy Grace Hopper, une des premières programmeuses du Harvard Mark I. Le compte précisait que Hopper serait basée sur une conception à module multi-puces, qui conduirait à un meilleur rendement de production et à des rebuts plus faibles^[17].

Lors du Nvidia GTC 2022, Nvidia a officiellement annoncé Hopper^[18].

• (en) Anne Elster et Tor Haugdahl, « Nvidia Hopper GPU and Grace CPU Highlights », Computing in Science & Engineering, vol. 24, n^o 2,‎ mars 2022, p. 95–100 (DOI 10.1109/MCSE.2022.3163817, Bibcode 2022CSE....24b..95E, S2CID 249474974, lire en ligne, consulté le 29 mai 2023)
• (en) Kohei Fujita, Takuma Yamaguchi, Yuma Kikuchi, Tsuyoshi Ichimura, Muneo Hori et Lalith Maddegedara, « Calculation of cross-correlation function accelerated by TensorFloat-32 Tensor Core operations on NVIDIA's Ampere and Hopper GPUs », Journal of Computational Science, vol. 68,‎ avril 2023 (DOI 10.1016/j.jocs.2023.101986 )
• (en) CUDA C++ Programming Guide, Nvidia, 17 avril 2023 (lire en ligne)
• (en) Hopper Tuning Guide, Nvidia, 13 avril 2023 (lire en ligne)
• (en) NVIDIA H100 Tensor Core GPU Architecture, Nvidia, 2022 (lire en ligne)

• (en) Jack Choquette, « NVIDIA Hopper H100 GPU: Scaling Performance », IEEE Micro, vol. 43, n^o 3,‎ mai 2023, p. 9–17 (DOI 10.1109/MM.2023.3256796, S2CID 257544490, lire en ligne, consulté le 29 mai 2023)
• (en) Samuel Moore, « Nvidia's Next GPU Shows That Transformers Are Transforming AI », IEEE Spectrum, 8 avril 2022 (consulté le 29 mai 2023)
• (en) Timothy Morgan, « Deep Dive Into Nvidia's "Hopper" GPU Architecture », The Next Platform, 31 mars 2022 (consulté le 29 mai 2023)

• Portail de l’informatique