Wie man AMD-GPUs mit „brrr“-Geschwindigkeit laufen lässt

• HipKittens ist eine Sammlung von Programmierprimitiven, die entwickelt wurde, um das Leistungspotenzial von AMD-GPUs auszuschöpfen, indem Speicherzugriffe, Scheduling und Cache-Wiederverwendung optimiert werden
• Die AMD-MI355X-GPU besitzt eine Struktur aus 256 Compute Units und 8 Chiplets (XCD) und bietet eine große Registerdatei sowie feingranulare Matrix-Core-Instruktionen
• Anders als NVIDIA verfügt AMD nicht über Register-Reallocation, asynchrone Matrix-Instruktionen oder mbarrier, weshalb statt Wave Specialization ein Scheduling mit 8-Wave Ping-Pong und 4-Wave Interleave effektiv ist
• HipKittens verbessert über chiplet-bewusstes (Grid-)Scheduling die Lokalität von L2- und LLC-Caches und erzielt bei GEMM- und Attention-Operationen höhere Bandbreite und TFLOPS
• Dieser Ansatz gleicht die mangelnde Software-Reife im AMD-GPU-Ökosystem aus und schafft eine Grundlage für bessere Skalierbarkeit von AI-Computing auf vielfältiger Hardware

---

Architektur und Leistungsmerkmale von AMD-CDNA-GPUs

• Die AMD-MI355X-GPU enthält 256 Compute Units (CU), wobei jede CU aus 4 SIMD-Einheiten besteht
  • Eine SIMD führt eine Wave aus, die aus 64 Threads besteht, im Gegensatz zu NVIDIAs Warp mit 32 Threads
• Die MI355X besitzt im Vergleich zur B200 nur etwa 70 % SRAM (165 KB) und bietet keine Funktionen wie asynchrone Matrix-Multiplikations-Instruktionen, Register-Reallocation, Tensor Memory Accelerator oder mbarrier
• Dafür bietet sie eine doppelt so große Registerdatei und 60 % mehr Prozessoren (256 CU gegenüber 160 SM)
  • Unterstützt werden kleine, feingranulare Matrix-Core-Instruktionen sowie direkte Global→Shared-Memory-Loads (ähnlich TMA)
• AMD setzt auf eine Chiplet-Architektur mit 8 Chiplets (XCD), wobei jedes XCD einen eigenen L2-Cache besitzt und darüber ein LLC-Cache liegt
• Laut Tabelle bietet die MI355X eine Rechenleistung von BF16 2,5 PFLOPs, MXFP8 5,0 PFLOPs, MXFP6 10,1 PFLOPs sowie 288 GB Speicherkapazität und 8 TB/s Bandbreite

Herausforderungen beim Kernel-Design für AMD

• Optimierung der Speicherzugriffe: Aufgrund von Einschränkungen des HIPCC-Compilers und nicht offengelegtem I/O-Verhalten ist das Design von Datenanordnung und Swizzle-Mustern entscheidend
• Scheduling innerhalb des Prozessors: AMD muss Registerdateien und kleine Matrix-Instruktionen statt Shared Memory nutzen
• Scheduling zwischen Prozessoren: Durch die Chiplet-basierte Struktur ist eine Aufgabenverteilung nötig, die NUMA-Effekte auf Cache-Ebene berücksichtigt

Speicherzugriffsmuster von HipKittens

• HipKittens (HK) verwendet Tiles als grundlegende Dateneinheit und bietet PyTorch-ähnliche Operationsfunktionen
  • Tiles werden durch Datentyp, Größe und Layout definiert und unterstützen über C++-Template-Metaprogrammierung vielfältige Eingaben
• Register-Scheduling: Da HIPCC bestimmte Register nicht als MFMA-Eingaben nutzen kann, bietet HK eine explizite Register-Fixierung
  • Entwickler können Register direkt festlegen und so Kernel mit maximaler Leistung schreiben
• Register-Layout: Bei AMD variiert das Layout je nach Datentyp und Matrixform, sodass kein einzelnes Swizzle-Muster ausreicht
  • Beispielsweise benötigen ein 16×16-bf16-Tile und ein 16×32-bf16-Tile unterschiedliche Swizzle-Muster
• Phasenstruktur von Instruktionen: AMDs Shared-Memory-Instruktionen haben nicht zusammenhängende Phasengruppen und sind intern nur unzureichend dokumentiert
  • HK liefert dafür einen per Reverse Engineering entwickelten Solver
• Adressgenerierung: AMD unterstützt asynchrone HBM→Shared-Memory-Loads und optimiert mit HBM-Adress-Swizzling

Scheduling innerhalb des Prozessors: Wave-Muster

• Wave Specialization ist bei NVIDIA effektiv, führt bei AMD wegen der fehlenden Register-Reallocation jedoch zu Leistungseinbußen
  • Producer-Waves belegen unnötige Register, während Consumer-Waves wegen Registermangels Spills verursachen
• Laut den Experimenten von HK führt Wave Specialization auf AMD zu geringerer arithmetischer Intensität und Speicherengpässen
  • Beispiel: Bei GEMM erreicht die HK-0/8-Konfiguration 1605 TFLOPs, CUTLASS 1570 TFLOPs
• Alternative Scheduling-Muster
  • 8-Wave Ping-Pong: Zwei Waves führen abwechselnd Speicher-/Compute-Cluster aus
  • 4-Wave Interleave: Eine Wave führt Speicher- und Compute-Schritte fein verzahnt abwechselnd aus
  • 8-Wave ist kompakter im Code, 4-Wave feingranularer, aber deutlich länger
  • Bei GEMM und Attention Forward erreicht 8-Wave Leistung auf SoTA-Niveau

Scheduling zwischen Prozessoren: chiplet-bewusster Ansatz

• Die AMD MI355X besitzt 8 XCD-Chiplets, von denen jedes über einen eigenständigen L2-Cache verfügt
  • Thread-Blöcke werden den Chiplets im Round-Robin-Verfahren zugewiesen, wodurch die Grid-Reihenfolge die Effizienz der Cache-Wiederverwendung direkt beeinflusst
• Ein einfaches Row-Major-Layout führt wegen geringer L2-Cache-Wiederverwendung zu Bandbreitenverlusten
  • Beispiel: L2 55 %, LLC 95 %, 15,1 TB/s, 1113 TFLOPs
• HK führt chiplet-bewusstes (Grid-)Scheduling ein und nutzt dabei gleichzeitig die Lokalität von L2- und LLC-Caches
  • Thread-Blöcke werden nach benachbarten Bereichen der Ausgabematrix gruppiert, um die Wiederverwendung von Eingabedaten zu maximieren

Praktische Kernel-Beispiele

• Die Hot Loops der Attention-Forward- und BF16-GEMM-Kernel verwenden den 8-Wave-Ping-Pong-Schedule von HK
  • Jede Schleife führt abwechselnd Compute–Memory-Cluster aus und synchronisiert über Schedule Barriers
  • Im Codebeispiel werden HK-Operationen wie mma_AtB, load, exp2, col_sum wiederholt verwendet

Fazit: AMD im Zeitalter von Multi-Silicon-AI

• HipKittens erreicht auf AMD CDNA3 und CDNA4 wettbewerbsfähige Leistung
  • Drei Kernelemente: optimierte Speicherzugriffe, AMD-zentriertes Wave-Scheduling und chiplet-bewusstes Grid-Scheduling
• HK-Kernel erzielen Spitzenleistung auf AMD und sind auch gegenüber NVIDIA-Blackwell-Kernels konkurrenzfähig
• Für mehr Vielfalt im AI-Computing ist ein breiterer Zugang zu AMD-GPUs nötig, und HipKittens bietet dafür eine zentrale Software-Grundlage
• Verbesserungen beim Register-Scheduling in AMDs HIPCC werden als wichtiger künftiger Entwicklungsbereich genannt

[GN⁺]

Hacker-News-Kommentare

• Empfehlenswert ist auch die Diskussion zu HipKittens
• Es gibt auch den Beitrag HipKittens: Fast and furious AMD kernels, der dieselbe Forschung behandelt. Dort gibt es Kommentare von George Hotz und AMD-Mitarbeitern
• Es ist erfreulich, dass sich die Wissenschaft mit solchen Problemen befasst, aber ich denke letztlich, dass AMD dieses Problem intern lösen muss
  • Ich finde, Hardware-Unternehmen sollten sich auf Hardware beschränken. So bleiben die Anreize sauber. Selbst wenn die Leistung dadurch um 20 % sinkt, halte ich das für besser
  • Stimme voll zu. AMD hat dieses Problem vor 10 Jahren aufgeschoben und versucht erst jetzt aufzuholen. Die Hardware ist großartig, aber wegen mangelnder Firmware-Kompetenz wird ihr Potenzial nicht ausgeschöpft
  • Dieses Forschungsteam hat allerdings auch ähnliche Software für Nvidia-GPUs entwickelt. Es sieht so aus, als würden exzellente Forscher einfach ihre Expertise ausspielen
  • Soweit ich weiß, arbeitet AMD bereits auf mehreren Ebenen an diesem Problem und kooperiert auch mit tinycorp
• Der Artikel vermittelt den Eindruck, dass die Optimierung bei AMD-GPUs wegen ihrer architektonischen Komplexität schwierig ist. Langfristig könnte AMDs Ansatz aber besser skalieren. Während Nvidia 2 Chiplets nutzt, hat AMD eine Struktur mit 8 Chiplets, wodurch Probleme mit der Speicherlokalität entstehen. In Zukunft wird die Zahl der Chiplets wohl weiter steigen, daher könnte die heutige Erfahrung im Umgang mit dieser Komplexität langfristig nützlich sein
  • AMD braucht für hohe Leistung keine Warp-Spezialisierung, was die Programmierung vereinfacht
• Viele Entwickler haben versucht, AMD-GPUs für normale Entwickler zum „go brrr“ zu bringen, sind aber gescheitert. Ich verstehe nicht, warum AMD die Softwareprobleme nicht selbst löst. Geld ist inzwischen genug da, und keine Entwickler einzustellen ist keine Ausrede mehr. Die GPUs fürs Rechenzentrum sind nicht schlecht, aber für eigene ML- und AI-Experimente ist Nvidia für Einzelpersonen immer noch deutlich besser. Mein fünf Jahre alter RTX 3090 fühlt sich immer noch besser an als jede bisher erschienene Consumer-GPU von AMD
  • Die Developer Experience bei AMD ist furchtbar. Sie nehmen nicht einmal Bug-Reports zu Treiberabstürzen an
  • Ich habe meinen Inferenz-Server kürzlich von einer NVidia 5090 auf zwei AMD R9700 32GB umgestellt, und es war durchweg eine positive Erfahrung. Unter dem Fedora-Kernel lief alles sofort ohne DKMS-Konfiguration, und mit ROCm war auch die Container-Anbindung einfach. Ich musste nur die Konfiguration von Ollama und Storyteller ändern. Es war eine deutlich angenehmere Erfahrung als mit CUDA
  • Nvidia pflegt sogar selbst einen Unreal-Engine-Fork. AMD spielt da nicht einmal annähernd in derselben Liga
  • Nvidia ist unter den Hardware-Unternehmen das einzige, das Software-Ingenieuren eine wettbewerbsfähige Vergütung bietet. Bei AMD gibt es noch immer eine Kultur, in der Software nicht als „richtige Arbeit“ gilt, und solche Trägheit lässt sich nur schwer ändern
• Mojo hatte Ideen, wie sich die Developer Experience (devX) auf AMD-GPUs verbessern ließe; mich würde interessieren, wie weit das gekommen ist
• Ich verstehe nicht, warum AMD nicht Milliarden Dollar in Softwareverbesserungen investiert. Nvidia ist das wertvollste Unternehmen der Welt, und AMD ist der einzige Konkurrent
  • AMD bemüht sich zwar, aber ich denke, es ist schwer, eine Organisationskultur, die jedes Jahr Hardware erneuert, in eine softwarezentrierte Kultur zu verwandeln. Software bringt nicht so unmittelbar Umsatz wie Hardware, daher neigt das Management dazu, sie niedriger zu priorisieren. Außerdem liefern externe Vendoren Code als Open Source, was kurzfristig gut aussieht, der langfristigen Qualität aber schadet. Wenn man einen Hardware-Trend auch nur einmal verpasst, riskiert man, hinter die Konkurrenz zurückzufallen
  • Ich habe bei mehreren GPU-Herstellern gearbeitet, und nur Nvidia betrachtet Software als Asset und investiert entsprechend. Andere Unternehmen sehen darin nur Kosten
• Ich persönlich mag das „go brr“-Meme nicht besonders, aber es ist witzig zu sehen, dass es an Orten wie Stanford verwendet wird
  • Tatsächlich wurde „go brr“ schon vor einem Jahr bei der Vorstellung von ThunderKittens verwendet
  • Wenn solche Memes auf offiziellen Kanälen von Universitäten auftauchen, ist das vielleicht schon ein Zeichen dafür, dass der Trend vorbei ist
• Das Projekt selbst ist großartig, aber ich frage mich, warum AMD so etwas nicht direkt selbst macht. AMD scheint die Bedeutung eines ausgereiften Software-Stacks immer noch nicht zu verstehen. Es braucht einen vereinheitlichten Stack wie CUDA, der auf allen Karten funktioniert. Früher glaubte ich, AMD würde irgendwann aufholen, aber inzwischen habe ich das fast aufgegeben
• Das Projekt ist gut, aber der Artikel selbst wirkt seltsam geschrieben
  • Der Text ist sehr unbeholfen. Entweder wurde zu stark auf AI gesetzt, oder jemand hat einen AI-Stil imitiert. Formulierungen wie „siehe part one“ oder „wie man AMD-GPUs go brr macht“ werden ständig wiederholt. Besonders schade ist, dass technische Inhalte, die man mit Grafiken hätte erklären sollen, stattdessen in 100 Zeilen Code ausgeschrieben werden