KI-gesteuerte Technologien verweben sich in unseren Alltag und haben das Potenzial, unseren Zugang zu Wissen zu verbessern und unsere Gesamtproduktivität zu steigern. Das Rückgrat dieser Anwendungen sind große Sprachmodelle (LLMs). LLMs sind speicherintensiv und erfordern in der Regel spezielle Hardwarebeschleuniger, um effizient bereitzustellen Dutzende Exaflops Rechenleistung. Dieser Blogbeitrag zeigt, wie wir die Rechenherausforderungen bewältigen können, indem wir den Speicher effektiver nutzen.
Der Großteil des Speichers und der Rechenleistung eines LLM wird verbraucht durch Gewichte In Matrix-Multiplikation Operationen. Durch die Verwendung schmalerer Datentypen reduziert den Speicherverbrauch. Beispielsweise kann das Speichern von Gewichten im 8-Bit-Format ganze Zahl (d. h. U8 oder S8) reduziert den Speicherbedarf um das Vierfache im Vergleich zu mit einfacher Genauigkeit (F32) und 2× relativ zu Halbpräzision (F16) oder bfloat16 (BF16). Außerdem vorherige Arbeit hat gezeigt, dass LLM-Modelle Matrixmultiplikationen mit Gewichte in S8 und Eingang in F16 (bei Beibehaltung einer höheren Genauigkeit der Benutzereingaben) ist eine effektive Methode zur Steigerung der Effizienz bei akzeptablen Kompromissen bei der Genauigkeit. Diese Technik ist bekannt als Nur-Gewicht-Quantisierung und erfordert eine effiziente Implementierung der Matrixmultiplikation mit gemischte Eingabenz. B. halbgenaue Eingabe multipliziert mit 8-Bit-Ganzzahl. Hardwarebeschleuniger, einschließlich GPUs, unterstützen einen festen Satz von Datentypen, und daher erfordert die Matrixmultiplikation mit gemischten Eingaben Softwaretransformationen, um sie den Hardwareoperationen zuzuordnen.
Zu diesem Zweck konzentrieren wir uns in diesem Weblog auf die Abbildung der Matrixmultiplikation mit gemischten Eingaben auf die NVIDIA Ampere-Architektur. Wir präsentieren Softwaretechniken, die sich mit der Datentypkonvertierung und Layoutkonformität befassen, um die Matrixmultiplikation mit gemischten Eingaben effizient auf hardwareunterstützte Datentypen und Layouts abzubilden. Unsere Ergebnisse zeigen, dass der Aufwand für zusätzliche Softwarearbeit minimal ist und eine Leistung nahe der maximalen Hardwarekapazität ermöglicht. Die hier beschriebenen Softwaretechniken werden im Open-Supply- NVIDIA/CUTLASS Repository.
 |
| Speicherbedarf für ein LLM-Modell mit 175 B Parametern und verschiedenen Datentypformaten. |
Die Matrix-Multiplikations-Akkumulations-Operation
Moderne KI-Hardwarebeschleuniger wie Googles TPU Und NVIDIAs GPU Multiplizieren Sie Matrizen nativ in der {Hardware}, indem Sie Tensor-Cores ansprechen, die spezialisierte Verarbeitungselemente sind, um Matrixoperationen zu beschleunigen, insbesondere für KI-Workloads. In diesem Weblog konzentrieren wir uns auf NVIDIA Ampere Tensor-Cores, die die Matrix-Multiplizieren-Akkumulieren (mma) Operation. Für den Relaxation des Blogs ist der Verweis auf mma ist für Ampere Tensor Cores. Die unterstützten Datentypen, Formen und das Datenlayout der beiden Eingangsmatrizen (Operanden genannt) für die mma Die Operationen sind in der {Hardware} festgelegt. Das bedeutet, dass Matrixmultiplikationen mit verschiedenen Datentypen und größeren Formen in der Software program implementiert werden, indem das Downside auf hardwareunterstützte Datentypen, Formen und Layouts aufgeteilt wird.
Der Tensor-Kern mma Operation wird durch die Angabe von zwei Eingangsmatrizen definiert (z. B. A und Bunten dargestellt), um eine Ergebnismatrix zu erzeugen, C. Der mma Der Vorgang unterstützt nativ gemischte Genauigkeit. Tensorkerne mit gemischter Genauigkeit Mischen von Eingaben zulassen (A Und B) Datentyp mit dem Ergebnis (C) Datentyp. Im Gegensatz dazu gemischte Eingabe Bei der Matrizenmultiplikation handelt es sich um ein Mischen der Eingabedatentypen. Da dies von der {Hardware} nicht unterstützt wird, muss es in der Software program implementiert werden.
 |
| Die Tensor Core-Operation M-mal-N-mal-Okay auf der Eingabematrix A von M-mal-Okay und der Matrix B von Okay-mal-N erzeugt die Ausgabematrix C von M-mal-N. |
Herausforderungen der Matrixmultiplikation mit gemischten Eingaben
Um die Diskussion zu vereinfachen, beschränken wir uns auf ein spezifisches Beispiel der Matrixmultiplikation mit gemischten Eingaben: F16 für die Benutzereingabe und U8 für die Modellgewichte (geschrieben als F16 * U8). Die hier beschriebenen Techniken funktionieren für verschiedene Kombinationen von Datentypen mit gemischten Eingaben.
Ein GPU-Programmierer hat Zugriff auf eine Hierarchie des Gedächtnisseseinschließlich globalem Speicher, gemeinsam genutztem Speicher und Registern, die in der Reihenfolge abnehmender Kapazität, aber zunehmender Geschwindigkeit angeordnet sind. NVIDIA Ampere Tensor Core mma Operationen verbrauchen Eingabematrizen aus Registern. Darüber hinaus müssen Eingabe- und Ausgabematrizen einem Datenlayout innerhalb einer Gruppe von 32 Threads entsprechen, das als KetteDer unterstützte Datentyp Und Format innerhalb eines Warps sind für eine mma Um eine Multiplikation mit gemischten Eingaben effizient zu implementieren, müssen daher die Herausforderungen der Datentypkonvertierung und Layoutkonformität in der Software program gelöst werden.
Datentypkonvertierung
Der mma Operation erfordert zwei Eingangsmatrizen mit demselben Datentyp. Daher erfordert die Matrixmultiplikation mit gemischten Eingängen, bei der einer der Operanden in U8 im globalen Speicher und der andere in F16 gespeichert ist, eine Datentypkonvertierung von U8 nach F16. Die Konvertierung bringt zwei Operanden nach F16 und bildet die gemischte Eingabe Matrixmultiplikation zur hardwaregestützten gemischte Genauigkeit Tensor-Kerne. Angesichts der großen Anzahl von Gewichten gibt es eine große Anzahl solcher Operationen, und unsere Techniken zeigen, wie ihre Latenz reduziert und die Leistung verbessert werden kann.
Layoutkonformität
Der mma Der Betrieb erfordert auch, dass das Format von zwei Eingangsmatrizen innerhalb der Register eines Warps mit der Hardwarespezifikation übereinstimmt. Das Format für die Eingangsmatrix B des U8-Datentyps bei der Matrixmultiplikation mit gemischten Eingängen (F16 * U8) muss mit dem konvertierten F16-Datentyp übereinstimmen. Dies wird als Layoutkonformität und muss in der Software program erreicht werden.
Die folgende Abbildung zeigt eine mma Operation verbrauchende Matrix A und Matrix B aus Registern zur Erzeugung einer Matrix C in Registern, verteilt auf einen Kettfaden. Der Faden T0 wird hervorgehoben und vergrößert, um die Gewichtsmatrix anzuzeigen B durchläuft eine Datentypkonvertierung und benötigt eine Layoutkonformität, um der hardwareunterstützten Tensor Core-Operation zugeordnet werden zu können.
Softwarestrategien zur Bewältigung von Herausforderungen
Eine typische Datentypkonvertierung umfasst eine Folge von Operationen an 32-Bit-Registern, wie unten dargestellt. Jeder rechteckige Block stellt ein Register dar und der angrenzende Textual content sind die Operationen. Die gesamte Sequenz zeigt die Konvertierung von 4xU8 nach 2x(2xF16). Die Sequenz umfasst ungefähr 10 Operationen.
Es gibt viele Möglichkeiten, Layoutkonformität zu erreichen. Zwei der bestehenden Lösungen sind:
- Gemeinsam genutzte Speicherlasten mit geringerer Bitbreite: Bei diesem Ansatz geben Threads Speicherladungen mit schmaler Bitbreite aus, die die U8-Daten vom gemeinsamen Speicher in Register verschieben. Dies führt zu zwei 32-Bit-Register, wobei jedes Register 2xF16 Werte enthält (siehe oben für die Matrix B’s Thread T0). Die schmalere Shared-Reminiscence-Final erreicht die Format-Konformität direkt in den Registern, ohne dass ein Shuffle erforderlich ist; allerdings wird dabei nicht die gesamte Bandbreite des Shared Reminiscence genutzt.
- Vorverarbeitung im globalen Speicher: Ein Alternativstrategie beinhaltet die Neuanordnung der Daten im globalen Speicher (eine Ebene über dem gemeinsam genutzten Speicher in Speicherhierarchie), wodurch breitere Shared-Reminiscence-Ladungen möglich werden. Dieser Ansatz maximiert die Bandbreitennutzung des Shared Reminiscence und stellt sicher, dass die Daten in einem konformen Format direkt in die Register geladen werden. Obwohl der Neuanordnungsprozess offline vor der LLM-Bereitstellung ausgeführt werden kann, sodass die Anwendungsleistung nicht beeinträchtigt wird, führt er einen zusätzlichen, nicht trivialen hardwarespezifischen Vorverarbeitungsschritt ein, der ein zusätzliches Programm zur Neuanordnung der Daten erfordert. NVIDIA/FasterTransformer wendet diese Methode an, um Herausforderungen hinsichtlich der Layoutkonformität wirksam zu bewältigen.
Optimierte Softwarestrategien
Um den Aufwand für die Datentypkonvertierung und Layoutkonformität weiter zu optimieren und zu reduzieren, haben wir FastNumericArrayConvertor Und FragmentShufflerjeweils.
FastNumericArrayConvertor operiert mit 4xU8 in 32-Bit-Registern, ohne einzelne 1xU8-Werte auszupacken. Darüber hinaus verwendet es weniger aufwendige Rechenoperationen, was die Anzahl der Anweisungen verringert und die Konvertierungsgeschwindigkeit erhöht.
Die Konvertierungssequenz für U8 bis F16 ist unten dargestellt. Die Operationen verwenden gepackte 32b-Register und vermeiden explizites Entpacken und Packen. FastNumericArrayConvertor verwendet die permute byte um Bytes von 4xU8 in zwei Register umzuordnen. Zusätzlich FastNumericArrayConvertor verwendet keine teuren Integer-zu-Fließkomma-Konvertierungsanweisungen und verwendet vektorisierte Operationen, um die gepackten Ergebnisse in zwei 32-Bit-Register mit 2x(2xF16) Werten. Die FastNumericArrayConvertor für U8 bis F16 werden ungefähr sechs Operationen verwendet, eine 1,6-fache Reduzierung im Vergleich zum oben gezeigten Ansatz.
 |
FastNumericArrayConvertor nutzt permute bytes und gepackte Arithmetik, wodurch die Anzahl der Anweisungen bei der Datentypkonvertierung reduziert wird. |
FragmentShuffler Behandelt die Layoutkonformität, indem Daten auf eine Weise neu gemischt werden, die die Verwendung von Ladevorgängen mit größerer Bitbreite ermöglicht, die Bandbreitennutzung des gemeinsam genutzten Speichers erhöht und die Gesamtzahl der Vorgänge verringert.
Die NVIDIA Ampere-Architektur bietet eine Ladematrixanweisung (ldmatrix). Der ldmatrix ist eine Operation auf Warp-Ebene, bei der 32 Threads eines Warps die Daten vom gemeinsamen Speicher in Register im Type Und Format Das mma Matrix A Und B konsumieren. Die Verwendung von ldmatrix reduziert die Anzahl der Ladeanweisungen und erhöht sich die Speicherbandbreitenauslastung. Da die ldmatrix Der Befehl verschiebt U8-Daten in Register, das Format nach dem Laden entspricht U8*U8 mma Operation, und nicht mit F16*F16 mma Betrieb. Wir haben FragmentShuffler um die Daten in den Registern mit shuffle neu anzuordnen (shfl.sync) Operationen zum Erreichen der Layoutkonformität.
Der wichtigste Beitrag dieser Arbeit besteht darin, Layoutkonformität durch Register-Shuffles zu erreichen, wodurch Offline-Vorverarbeitung im globalen Speicher oder Shared-Reminiscence-Ladungen mit geringerer Bitbreite vermieden werden. Darüber hinaus bieten wir Implementierungen für FastNumericArrayConvertor umfasst die Datentypkonvertierung von U8 bis F16, S8 bis F16, U8-bis-BF16Und S8 bis BF16.
Leistungsergebnisse
Wir haben die Leistung von acht Blended-Enter-Varianten von Unsere Methode (unten in blau und rot dargestellt; Variation der Datentypen der Matrix A Und B) und zwei gemischte Genauigkeit Datentypen (grün dargestellt) auf einem NVIDIA A100 SXM-Chip. Die Leistungsergebnisse sind in FLOPS (höher ist besser). Insbesondere erfordern die ersten acht Matrixmultiplikationen im Vergleich zu den letzten beiden zusätzliche Operationen, da die Varianten mit gemischter Genauigkeit direkt auf hardwarebeschleunigte Tensor Core-Operationen abzielen und keine Datentypkonvertierung und Layoutkonformität benötigen. Trotzdem zeigt unser Ansatz eine Leistung bei Matrixmultiplikationen mit gemischten Eingaben, die nur geringfügig unter oder auf Augenhöhe mit der gemischten Genauigkeit liegt.
 |
Leistung der Matrixmultiplikation mit gemischten Eingaben auf dem NVIDIA A100 40 GB SMX4-Chip für eine rechenintensive Matrixproblemform m=3456, n=4096, ok=2048. |
Danksagung
Wir möchten mehrere Personen erwähnen, die durch technisches Brainstorming und die Verbesserung des Blogbeitrags beigetragen haben, darunter Quentin Colombet, Jacques Pienaar, Allie Culp, Calin Cascaval, Ashish Gondimalla, Matt Walsh, Marek Kolodziej und Aman Bhatia. Wir möchten unseren NVIDIA-Partnern Rawn Henry, Pradeep Ramani, Vijay Thakkar, Haicheng Wu, Andrew Kerr, Matthew Properly und Vartika Singh danken.