Optimisations de performances de multiplication de la multiplication à la matrice à entrée mixte

Les technologies axées sur l’IA se tient dans le tissu de nos routines quotidiennes, avec le potentiel d’améliorer notre accès aux connaissances et de stimuler notre productivité globale. L’épine dorsale de ces applications réside dans les modèles de grande langue (LLMS). Les LLM sont à forte intensité de mémoire et nécessitent généralement des accélérateurs matériels spécialisés pour livrer efficacement des dizaines d’exaflops de la puissance de calcul. Cet article de blog montre comment nous pouvons commencer à relever les défis de calcul en utilisant la mémoire plus efficacement.

La majeure partie de la mémoire et du calcul d’un LLM sont consommés par poids dans Multiplication matricielle opérations. Utiliser plus étroit types de données réduit la consommation de mémoire. Par exemple, stocker des poids dans le 8 bits entier (IE, U8 ou S8) Le type de données réduit l’empreinte de la mémoire par 4 × par rapport à à une seule précision (F32) et 2 × par rapport à demi-précision (F16) ou bfloat16 (BF16). En outre, Les travaux antérieurs ont a montré que les modèles LLM exécutant des multiplications matricielles avec poids Dans S8 et saisir Dans F16 (préserver une précision plus élevée de l’entrée utilisateur) est une méthode efficace pour augmenter l’efficacité avec des compromis acceptables en précision. Cette technique est connue sous le nom quantification de poids uniquement et nécessite une implémentation efficace de la multiplication matricielle avec mixtespar exemple, l’entrée de demi-précision multipliée par un entier 8 bits. Les accélérateurs matériels, y compris les GPU, prennent en charge un ensemble fixe de types de données et, par conséquent, la multiplication matricielle à entrée mixte nécessite des transformations de logiciels pour mapper les opérations matérielles.

À cette fin, dans ce blog, nous nous concentrons sur la cartographie de la multiplication matricielle à entrée mixte sur le Architecture nvidia ampère. Nous présentons des techniques logicielles traitant efficacement la conversion du type de données et la conformité de la mise en place de la multiplication de la matrice à entrée mixte sur les types de données et les dispositions soutenus par le matériel. Nos résultats montrent que les frais généraux de travail supplémentaire dans les logiciels sont minimes et permet des performances proches des capacités matérielles maximales. Les techniques logicielles décrites ici sont publiées dans l’ouverture open Nvidia / cutlass dépôt.

Empreinte mémoire pour un modèle LLM paramètre 175B avec divers formats de types de données.

L’opération matricielle-accumulée-accumulée

Accélérateurs de matériel AI modernes tels que TPU de Google et GPU de Nvidia Multipliez les matrices nativement dans le matériel en ciblant les noyaux du tenseur, qui sont des éléments de traitement spécialisés pour accélérer les opérations matricielles, en particulier pour les charges de travail de l’IA. Dans ce blog, nous nous concentrons sur les noyaux de tenseur nvidia ampère, qui fournissent le matrice-accumulée (mma) opération. Pour le reste du blog, la référence à mma est pour les noyaux de tenseur ampère. Les types de données pris en charge, les formes et la disposition des données des deux matrices d’entrée (appelées opérandes) pour le mma L’opération est fixée dans le matériel. Cela signifie que les multiplications matricielles avec divers types de données et des formes plus grandes sont implémentées dans le logiciel en malissant le problème sur des types de données, des formes et des dispositions soutenues par le matériel.

Le noyau du tenseur mma Le fonctionnement est défini en spécifiant deux matrices d’entrée (par exemple, UN & Billustré ci-dessous) pour produire une matrice de résultats, C. Le mma L’opération soutient nativement la précision mixte. Noyaux de tenseur de précision mixte Autoriser l’entrée de mélange (UN et B) Type de données avec le résultat (C) Type de données. En revanche, mixte La multiplication matricielle implique le mélange des types de données d’entrée, et il n’est pas pris en charge par le matériel, il doit donc être implémenté dans le logiciel.

Le fonctionnement du cœur du tenseur de m-by-n-by-k sur la matrice d’entrée A de m-by-k et la matrice B de K-by-N produit la matrice de sortie C de m-by-n.

Défis de la multiplication matricielle à entrée mixte

Pour simplifier la discussion, nous nous restreignions à un exemple spécifique de multiplication matricielle à entrée mixte: F16 pour l’entrée de l’utilisateur et U8 pour les poids du modèle (écrit comme F16 * U8). Les techniques décrites ici fonctionnent pour diverses combinaisons de types de données à entrée mixte.

Un programmeur GPU peut accéder à un Hiérarchie de mémoirey compris la mémoire globale, la mémoire partagée et les registres, qui sont organisés par ordre diminution de la capacité mais augmentant la vitesse. Nvidia ampère noyau tenseur mma Les opérations consomment des matrices d’entrée des registres. De plus, des matrices d’entrée et de sortie sont nécessaires pour se conformer à une disposition des données dans un groupe de 32 threads appelés un chaîne. Le type de données pris en charge et La disposition dans une chaîne est fixée pour un mma Fonctionnement, pour mettre en œuvre efficacement la multiplication à entrée mixte, il est nécessaire de résoudre les défis de la conversion du type de données et de la conformité de la disposition dans le logiciel.

Conversion de type de données

Le mma L’opération nécessite deux matrices d’entrée avec le même type de données. Ainsi, la multiplication matricielle à entrée mixte, où l’un des opérandes est stocké en U8 dans la mémoire globale et autres en F16, nécessite une conversion de type de données de U8 en F16. La conversion apportera deux opérandes à F16, cartographiant le mixte Multiplication matricielle à l’approbation du matériel précision Noyaux du tenseur. Compte tenu du grand nombre de poids, il existe un grand nombre de ces opérations et nos techniques montrent comment réduire leur latence et améliorer les performances.

Conformité de mise en page

Le mma L’opération nécessite également que la disposition de deux matrices d’entrée, dans les registres d’une chaîne, soit conforme à la spécification matérielle. La disposition de la matrice d’entrée B du type de données U8 dans la multiplication matricielle à entrée mixte (F16 * U8) doit se conformer au type de données F16 converti. C’est ce qu’on appelle conformité de mise en page et doit être réalisé dans le logiciel.

La figure ci-dessous montre un mma opération consommatrice de matrice UN et matrice B des registres à produire une matrice C dans les registres, distribué sur une chaîne. Le fil T0 est mis en surbrillance et zoom pour montrer la matrice de poids B passe par la conversion du type de données et a besoin d’une conformité de mise en page pour pouvoir se glisser vers le fonctionnement du noyau du tenseur soutenu par le matériel.

Stratégies logicielles relève des défis

Une conversion de type de données typique implique une séquence d’opérations sur des registres 32 bits, illustrés ci-dessous. Chaque bloc rectangulaire représente un registre et le texte adjacent est les opérations. La séquence entière montre la conversion de 4XU8 à 2x (2xf16). La séquence implique environ 10 opérations.

Il existe de nombreuses façons d’atteindre la conformité de mise en page. Deux des solutions existantes sont:

1. Bitwidth plus étroit Charges de mémoire partagée: Dans cette approche, les threads émettent des charges de mémoire BitWidth étroites en déplaçant les données U8 de la mémoire partagée aux registres. Il en résulte deux Registres 32 bits, avec chaque registre contenant des valeurs 2xf16 (illustré ci-dessus pour la matrice BFil T0). La charge de mémoire partagée plus étroite atteint la conformité de mise en page directement dans les registres sans avoir besoin de mélanges; Cependant, il n’utilise pas la bande passante de mémoire partagée complète.
2. Prétraitement dans la mémoire globale: Un stratégie alternative implique le réorganisation des données dans la mémoire globale (un niveau supérieur à la mémoire partagée en hiérarchie de mémoire), permettant des charges de mémoire partagées plus larges. Cette approche maximise l’utilisation de la bande passante de mémoire partagée et garantit que les données sont chargées dans une disposition conforme directement dans les registres. Bien que le processus de réarrangement puisse être exécuté hors ligne avant le déploiement LLM, n’assurant aucun impact sur les performances de l’application, il introduit une étape de prétraitement spécifique au matériel supplémentaire et non triviale qui nécessite un programme supplémentaire pour réorganiser les données. Nvidia / FasterTransformateur adopte cette méthode pour relever efficacement les défis de conformité de mise en page.

Stratégies logicielles optimisées

Pour optimiser et réduire davantage les frais généraux de la conversion du type de données et de la conformité de la disposition, nous avons mis en œuvre FastNumericArrayConvertor et FragmentShufflerrespectivement.

FastNumericArrayConvertor fonctionne sur 4XU8 dans des registres 32 bits sans déballer les valeurs individuelles 1xu8. De plus, il utilise des opérations arithmétiques moins coûteuses qui réduisent le nombre d’instructions et augmentent la vitesse de la conversion.

La séquence de conversion pour U8 à F16 est indiqué ci-dessous. Les opérations utilisent des registres 32b emballés, en évitant le déballage et l’emballage explicites. FastNumericArrayConvertor utilise le permute byte Pour réorganiser les octets de 4XU8 en deux registres. En plus, FastNumericArrayConvertor n’utilise pas en entier coûteux pour les instructions de conversion à point flottante et utilise des opérations vectorisées pour obtenir les résultats emballés deux Registres 32 bits contenant des valeurs 2x (2xf16). Le FastNumericArrayConvertor Pour U8 à F16, utilise approximativement six opérations, une réduction de 1,6 × par rapport à l’approche indiquée ci-dessus.

FastNumericArrayConvertor utilise permute bytes et arithmétique emballée, réduisant le nombre d’instructions dans la conversion du type de données.

FragmentShuffler gère la conformité de la disposition en mélangeant les données d’une manière qui permet d’utiliser un fonctionnement plus large de charge Bitwidth, d’augmenter l’utilisation de la bande passante de mémoire partagée et de réduire le nombre total d’opérations.

L’architecture Nvidia Ampère fournit une instruction de matrice de charge (ldmatrix). Le ldmatrix est une opération au niveau de la chaîne, où 32 threads d’une déformation déplacent les données de la mémoire partagée vers les registres dans le forme et mise en page que mma matrice UN et B consommer. L’utilisation de ldmatrix réduit le nombre d’instructions de chargement et augmentation L’utilisation de la bande passante de la mémoire. Depuis le ldmatrix L’instruction déplace les données U8 vers les registres, la disposition après la charge se conforme à U8 * U8 mma Fonctionnement, et non avec F16 * F16 mma opération. Nous avons mis en œuvre FragmentShuffler Pour réorganiser les données dans les registres à l’aide de Shuffle (shfl.sync) opérations pour réaliser la conformité de la disposition.

La contribution la plus significative de ce travail consiste à atteindre la conformité de la disposition par le biais de shuffles de registre, d’éviter le prétraitement hors ligne dans la mémoire globale ou les charges de mémoire partagées BitWidth plus étroites. De plus, nous fournissons des implémentations pour FastNumericArrayConvertor couvrant la conversion de type de données à partir de U8 à F16, S8 à F16, U8 à BF16et S8 à BF16.

Résultats de performance

Nous avons mesuré les performances de huit variantes à entrée mixte de Notre méthode (montré ci-dessous en bleu et rouge; variant les types de données de matrice UN et B) et deux précision Types de données (illustrés en vert) sur une puce NVIDIA A100 SXM. Les résultats des performances sont présentés dans Flops (plus haut, c’est mieux). Notamment, les huit premières multipies matricielles nécessitent des opérations supplémentaires par rapport aux deux derniers, car les variantes de précision mixte ciblent directement les opérations de noyau de tenseur accéléré par le matériel et n’ont pas besoin de conversion de type de données et de conformité de mise en page. Même ainsi, notre approche démontre les performances de multiplication de la matrice à entrée mixte que légèrement en dessous ou à égalité avec une précision mixte.

Performances de multiplication matricielle à entrée mixte sur la puce NVIDIA A100 40GB SMX4 pour une forme de problème de matrice liée à la calcul m=3456, n=4096, k=2048.

Remerciements

Nous aimerions mentionner plusieurs personnes qui ont contribué à un brainstorming technique et à l’amélioration du blog, notamment Quentin Colombet, Jacques Pienaar, Allie Culp, Calin Cascaval, Ashish Gondimalla, Matt Walsh, Marek Kolodziej et Aman Bhatia. Nous tenons à remercier nos partenaires Nvidia Rawn Henry, Pradeep Ramani, Vijay Thakkar, Haicheng Wu, Andrew Kerr, Matthew bien et Vartika Singh.