Merkmale der H100 GPU
- Bietet 80 GB HBM3-Speicher und 3 TB/s Bandbreite (tatsächlich etwas niedriger)
- Bietet 50 MB L2-Cache und 12 TB/s Bandbreite. Auf der GPU in zwei 25-MB-Sektionen aufgeteilt und per Crossbar verbunden (die Crossbar ist ein Faktor für Leistungseinbußen)
- Besteht aus 132 Streaming Multiprocessors (SM), wobei jeder SM wie folgt aufgebaut ist:
- Bietet bis zu 227 KB Shared Memory innerhalb eines 256-KB-L1-Caches (zusammen etwa 33 TB/s Bandbreite)
- Bietet einen Tensor Memory Accelerator (TMA), der asynchrone Adressgenerierung und Memory Fetching ermöglicht. Unterstützt auch Funktionen wie ein On-Chip-Memory-Netzwerk, die in diesem Beitrag jedoch nicht behandelt werden
- In 4 Quadranten unterteilt, wobei jeder Quadrant aus einem Warp Scheduler, 512 Vector Registers (jeweils mit 32 4-Byte-Wörtern), einem Tensor Core für Matrixmultiplikation und integrierten Instructions für parallele Operationen wie Summe/Multiplikation besteht
Tipps zur Leistungsoptimierung der H100 GPU
- Es ist wichtig, die Tensor Cores maximal auszunutzen. Die H100 erreicht 989 TFLOPs FP16-Matrixmultiplikationsleistung, daher hängt die Gesamtauslastung der GPU stark davon ab, wie gut die Tensor Cores genutzt werden
- Um die Tensor Cores maximal auszunutzen, müssen jedoch die folgenden Punkte berücksichtigt werden:
- Die Nutzung der Warp Group Matrix Multiply Accumulate (WGMMA)-Instructions ist unverzichtbar, aber schwierig in der Anwendung
- Shared Memory ist nicht so schnell wie erwartet und erfordert bei der Nutzung besondere Sorgfalt
- Adressgenerierung ist teuer, daher sollte mit Optimierungen wie dem Tensor Memory Accelerator (TMA) gearbeitet werden
- Eine höhere Occupancy ist weiterhin hilfreich, und Registers sind die wichtigste Ressource
- Diese Eigenschaften gelten bis zu einem gewissen Grad nicht nur für die H100, sondern auch für andere GPUs, allerdings ist die Nutzung der Tensor Cores bei der H100 besonders wichtig und anspruchsvoll
ThunderKittens: in H100 optimiertes eingebettetes CUDA-DSL
- Ein auf CUDA basierendes eingebettetes DSL, das entwickelt wurde, um die Leistung moderner GPUs wie der NVIDIA H100 maximal auszureizen
- Bietet die folgenden 4 tile-basierten Template-Typen:
- Register Tile (2D-Tensor, der in Registers gespeichert ist)
- Register Vector (1D-Tensor, der in Registers gespeichert ist)
- Shared Tile (2D-Tensor, der im Shared Memory gespeichert ist)
- Shared Vector (1D-Tensor, der im Shared Memory gespeichert ist)
- Bietet außerdem verschiedene Operatoren zur Tile-Manipulation (exp, mul, sum usw.) auf Warp- oder Warp-Group-Ebene
- Die Implementierung bestehender Flash Attention- und Flash Attention 2-Kernels mit ThunderKittens vereinfacht den Code deutlich und bringt auf der H100 bis zu 30 % mehr Leistung
- Auch der Based Linear Attention-Kernel wurde mit ThunderKittens implementiert und erreicht 215 TFLOPs Leistung (aufgrund der Algorithmus-Eigenschaften über 300 TFLOPs einschließlich Recompute)
Betrachtung aus philosophischer Perspektive
- ThunderKittens funktioniert so gut, weil es nicht versucht, alles abzudecken, sondern eine einfache tile-basierte Abstraktion bietet, die sehr gut zur GPU-Architektur passt
- Die Nutzung von 1024-Bit-Vector-Registers statt traditioneller 32-Bit-Wörter ist ein Fortschritt, aber es braucht auch einen Paradigmenwechsel hin dazu, ein 16x16-Tile als Einheit des Registers zu betrachten
- Da AI-Workloads letztlich vor allem aus Matrixmultiplikation, Reduction und Reshape bestehen, ist ein tile-basierter Ansatz plausibel; auch aus Hardware-Sicht wird sich die Entwicklung wohl über Systolic Arrays hinaus in Richtung Unterstützung kleiner Matrixmultiplikationen bewegen
- Darüber hinaus ist ein Umdenken nötig hin zu einer Gestaltung von AI-Algorithmen in hardwareoptimierter Form. So sollte etwa die Größe des Zustands eines RNN auf das Maß begrenzt werden, das in einen SM passt, und auch die Rechendichte sollte an das vom Hardware-Design geforderte Niveau angepasst werden
Meinung von GN⁺
- ThunderKittens kann für Entwickler, die mit CUDA vertraut sind, eine attraktive Option sein, weil sich Leistungsgewinne erzielen lassen, ohne bestehenden Kernel-Code grundlegend anzufassen
- Für Einsteiger kann die Einstiegshürde jedoch weiterhin hoch sein. Künftig wären mehr Beispielcode und Lernmaterialien wünschenswert
- Interessant ist der Plan, ThunderKittens-Unterstützung nicht nur auf H100, sondern auch auf andere GPUs wie AMD auszuweiten. Das könnte helfen, die Vendor-Abhängigkeit zu verringern
- Letztlich ist es ein sehr wichtiger Punkt, AI-Modelle und -Algorithmen selbst in hardwareoptimierter Form zu entwerfen. Dafür ist allerdings zunächst ein tiefes Verständnis der Hardware-Eigenschaften nötig, und ThunderKittens kann Entwicklern helfen, solche Einsichten zu gewinnen
- Es ist zu erwarten, dass die kontinuierliche Forschung und Open-Source-Arbeit von Hazy Research das CUDA-Ökosystem deutlich beleben wird. Langfristig scheint jedoch auch das Aufkommen von Frameworks mit höherem Abstraktionsniveau notwendig
1 Kommentare
Hacker-News-Kommentar
Die Anforderungen an AI-Hardware werden immer klarer. GPUs wurden ursprünglich für andere Zwecke entwickelt, werden aber für AI eingesetzt, weil sie gute Hardware für Matrixmultiplikation besitzen. Eine "AI-GPU" könnte auf einige Funktionen einer echten GPU verzichten, und der Trend geht zu kürzeren Zahlenformaten (16-Bit-, 8-Bit-, 2-Bit- und 1-Bit-Gleitkommazahlen). Dieses Paper deutet darauf hin, dass Hardware, die 16x16-Tiles bevorzugt, viele Vorteile bietet.
Es braucht AI-spezifische Co-Prozessoren (NPUs), besonders in prosumerorientierten Desktop-Systemen für Entwickler, Fachleute und Gamer. GPUs funktionieren in Unternehmen, sind aber für AI im Bereich Personal Computing unhandlich. Besonders problematisch sind die VRAM-Beschränkungen und das Fehlen einer standardisierten offenen API außer Vulkan.
Um die AI-Forschung voranzubringen, müssen Neurowissenschaften und Psychologie besser erforscht werden. Möglicherweise spielen auch Dinge eine Rolle, die mit der Graph-Topologie neuronaler Netze zusammenhängen.
Ist das so etwas wie eine benutzerfreundliche Version von CUTLASS?
Das ThunderKittens-Maskottchen hat etwas von einer Katze / Sony-Aibo-Ästhetik. Es sieht so aus, als wäre es gut mit AI generiert worden.
Wenn Universal Basic Compute (UBC) als Ersatz für Universal Basic Income in Betracht gezogen würde, wäre das eine ziemlich dystopische Zukunft. Man stelle sich vor, ein einziges Unternehmen wie Nvidia würde das gesamte Computing herstellen.