Die explosive Leistungssteigerung von GPUs

 (hazyresearch.stanford.edu)
1 Punkte von GN⁺ 2024-05-13 | 1 Kommentare | Auf WhatsApp teilen
  • Da die Kosten für AI-Berechnungen steigen, fasst Hazy Research zusammen, dass der Schlüssel zur Optimierung der GPU-Leistung darin liegt, die Tensor Cores der NVIDIA H100 ohne Unterbrechung auszulasten
  • Die H100 erreicht bei Half-Precision-Matrixmultiplikation 989 TFLOPs, kommt bei allgemeinen Operationen jedoch nur auf etwa 60 TFLOPs; sobald die Tensor Cores stillstehen, fällt die Auslastung stark ab
  • Um sich der Maximalleistung zu nähern, müssen WGMMA, Shared-Memory-Anordnung, Adressgenerierung und Occupancy gemeinsam behandelt werden; ohne wgmma.mma_async bleibt man in Mikrobenchmarks bei rund 63 % des Peaks
  • Das veröffentlichte eingebettete CUDA-DSL ThunderKittens kapselt mit Tile- und Vektor-Abstraktionen Komplexität wie Swizzling und Register-Layout und vereinfacht so das Schreiben von FlashAttention-artigen Kerneln
  • Der FlashAttention-2-Forward-Kernel für die H100 lässt sich in rund 100 Zeilen schreiben und ist etwa 30 % schneller als FlashAttention-2; der Based-Linear-Attention-Kernel läuft mit 215 TFLOPs

Bedingungen, die die H100-Leistung bestimmen

  • AI nutzt viel Compute, und Hazy Research hat in den vergangenen Jahren daran gearbeitet, dass AI entweder mit weniger Compute auskommt oder auf gegebener Compute effizienter läuft
  • Das praktische Ziel ist, die Erkenntnisse aus der Beschleunigung von GPUs zusammenzufassen und das eingebettete CUDA-DSL ThunderKittens zu veröffentlichen, das beim Schreiben schneller Kernel helfen soll
  • Allgemeiner geht es darum, wie das Verständnis von Hardware die Sicht auf AI-Compute verändert hat

Architektur und Engpässe der NVIDIA H100

  • Diskutiert wird auf Basis einer H100-SXM-GPU mit folgender Konfiguration
    • 80 GB HBM3, Bandbreite 3 TB/s
    • 50 MB L2-Cache, Bandbreite 12 TB/s, aufgeteilt in zwei 25-MB-Sektionen über die GPU hinweg und per Crossbar verbunden

  • 132 SMs

    • Jeder SM besitzt einen 256-KB-L1-Cache inklusive bis zu 227 KB Shared Memory und erreicht zusammen etwa 33 TB/s Bandbreite
    • Die neue Hopper-Hardware Tensor Memory Accelerator (TMA) übernimmt asynchrone Adressgenerierung und Memory-Fetches
    • Jeder SM besteht aus 4 Quadranten, und jeder Quadrant enthält einen Warp-Scheduler, 512 Vektorregister, Tensor Cores für Matrixmultiplikation sowie parallele Spezialinstruktionen
    • Sämtliche Compute-Arbeit findet im SM statt, der Großteil in den Registern
    • Der Schlüssel zu hoher Leistung auf der H100 ist, die Tensor Cores kontinuierlich mit Daten zu versorgen
    • Die H100 liefert 989 TFLOPs bei Half-Precision-Matrixmultiplikation und etwa 60 TFLOPs bei „allen anderen“ Operationen
    • In Zyklen, in denen Tensor Cores genutzt werden, wird mindestens 94 % Hardware-Auslastung erreicht
    • In Zyklen ohne Tensor-Core-Nutzung bleibt die Auslastung bei höchstens 6 %

WGMMA: notwendig, aber schwierig

  • Die H100 besitzt die warp group matrix multiply accumulate-Instruktion wgmma.mma_async
    • In PTX: wgmma.mma_async
    • In SASS: HGMMA / IGMMA / QGMMA / BGMMA
  • Auf früheren GPUs waren wmma.mma.sync und mma.sync synchron: Ein Warp aus 32 Threads fütterte die Tensor Cores mit Daten und wartete auf das Ergebnis
  • wgmma.mma_async synchronisiert stattdessen 128 zusammenhängende Threads, die über alle Quadranten eines SM kooperieren, und startet die Matrixmultiplikation asynchron direkt aus dem Shared Memory
    • Die Warps können während der laufenden Matrixmultiplikation andere Arbeiten in Registern ausführen
    • Auf die Ergebnisse kann zum gewünschten Zeitpunkt gewartet werden
  • In Mikrobenchmarks waren diese Instruktionen nötig, um die volle Compute-Leistung der H100 herauszuholen
    • Ohne sie blieb die GPU bei etwa 63 % der Spitzen-Auslastung
    • Ein möglicher Grund ist, dass Tensor Cores selbst bei lokalen Ressourcen tiefe Hardware-Pipelines benötigen
  • Die größte Schwierigkeit ist die Komplexität des Memory-Layouts
    • Ein unswizzeltes Shared-Memory-Layout hat sehr schlechte Coalescing-Eigenschaften und beansprucht viel L2-Bandbreite
    • Das swizzelte Layout war schwer zu verstehen, auch weil die Dokumentation fehlerhaft war
    • Das swizzelte Layout scheint nur für bestimmte Matrix-Shapes zu funktionieren und passt nicht gut zu anderen Funktionen von wgmma.mma_async
    • Die Hardware kann auf dem Weg zum Tensor Core zwar eine Submatrix transponieren, aber nur wenn das Layout nicht swizzelt ist
  • In Kerneln wie FlashAttention können TMA und L2-Cache dieses Problem teilweise kaschieren, weil sie schnell genug sind
  • Um die Hardware vollständig auszunutzen, müssen Memory-Requests gut coalesced sein und Bank Conflicts vermieden werden; deshalb ist die Kontrolle über das Layout wichtig

Shared Memory und Bank Conflicts

  • Die Single-Access-Latenz von Shared Memory liegt offenbar bei rund 30 Zyklen; in dieser Zeit könnten die Tensor Cores eines SM fast zwei 32x32-Matrixmultiplikationen ausführen
  • Frühere Arbeiten wie FlashAttention konzentrierten sich vor allem auf den HBM-SRAM-Engpass, und historisch war dieser Engpass auch tatsächlich wichtig
  • Da HBM schneller geworden ist und Tensor Cores stärker skaliert haben als andere Chipbereiche, müssen inzwischen selbst kleine Shared-Memory-Latenzen entfernt oder verborgen werden
  • Shared Memory ist in 32 Banks aufgeteilt, sodass ohne Vorsicht Bank Conflicts entstehen
    • Wenn mehrere unterschiedliche Speicherstücke gleichzeitig aus derselben Memory-Bank angefordert werden, werden die Requests serialisiert
    • Erfahrungsgemäß kann ein Kernel dadurch unverhältnismäßig stark verlangsamt werden
    • Die von WGMMA und MMA geforderten Register-Layouts können bei naiver Nutzung Bank Conflicts verursachen
  • Die Lösung besteht darin, Shared Memory mit verschiedenen Swizzling-Mustern neu anzuordnen, um Konflikte zu vermeiden
  • Wo möglich, sollte man Transfers zwischen Registern und Shared Memory vermeiden; wenn sie nötig sind, sind asynchrone Datenbewegungen über Spezialhardware wie WGMMA und TMA vorteilhaft
  • Tatsächliche synchrone Transfers mit normalen Warps sind zwar am allgemeinsten, kommen aber dem schlechtesten Fallback nahe

Adressgenerierung und TMA

  • Auf der H100 sind sowohl Tensor Cores als auch Memory so schnell, dass bereits die Erzeugung der zu ladenden Memory-Adressen einen erheblichen Teil der Chip-Ressourcen beansprucht
    • Bei komplexen Interleaving- oder Swizzling-Mustern fällt das noch stärker ins Gewicht
  • Der Tensor Memory Accelerator (TMA) von NVIDIA erlaubt es, multidimensionale Tensor-Layouts in Global und Shared Memory festzulegen, asynchron ein Subtile dieses Tensors zu laden und bei Abschluss eine Barrier auszulösen
  • TMA senkt die Kosten der Adressgenerierung und vereinfacht außerdem den Aufbau von Pipelines
  • TMA wird – ähnlich wie wgmma.mma_async – als essenziell angesehen, um das Potenzial der H100 auszuschöpfen
    • Nach praktischer Erfahrung ist TMA möglicherweise sogar wichtiger als WGMMA
    • Es spart Register-Ressourcen und Instruction-Dispatch
    • Es kann auch asynchrone Reduction auf Global Memory ausführen, was in komplexen Backward-Kerneln nützlich ist
  • Auch bei TMA war etwas Reverse Engineering nötig, um die Swizzling-Modi zu verstehen, aber es war weniger schmerzhaft als bei WGMMA

Occupancy verbirgt Kosten

  • Occupancy bezeichnet in CUDA die Zahl der Threads, die gleichzeitig auf derselben Ausführungshardware eingeplant sind
  • Der Warp-Scheduler eines SM-Quadranten versucht in jedem Zyklus, einer Warp eine Instruktion zuzuweisen, die dafür bereit ist
  • Die H100 ist in mancher Hinsicht weniger von Occupancy abhängig als frühere Generationen
    • Dank asynchroner Funktionen kann schon ein einzelner Instruction-Stream gleichzeitig Memory-Fetch, Matrixmultiplikation, Shared-Memory-Reduction und Register-Arithmetik auslasten
  • Dennoch ist Occupancy sehr nützlich, um Fehler und Synchronisationskosten zu verbergen
    • Eine perfekt entworfene Pipeline kann auch ohne zusätzliche Occupancy schnell sein
    • In der Praxis wirkte es jedoch so, als seien NVIDIA-GPUs mit Blick auf Occupancy entworfen worden
    • Da es viele Synchronisationspunkte und Fehlermöglichkeiten gibt, verbessert höhere Occupancy oft die tatsächlich erreichte Hardware-Auslastung
  • Auf der H100 ist Occupancy hilfreich, auf der A100 und der RTX 4090 offenbar jeweils noch wichtiger
    • Als möglicher Grund wird genannt, dass diese GPUs stärker auf synchronen Instruction-Dispatch angewiesen sind als die H100

ThunderKittens: ein kleines DSL in CUDA

  • ThunderKittens ist ein eingebettetes CUDA-DSL, das entwickelt wurde, um auf der H100 schnellere Kernel einfacher schreiben zu können
  • Ursprünglich entstand es für den internen Einsatz im Labor und wurde später veröffentlicht
  • Der Name wurde gewählt, weil Kätzchen niedlich sind und weil es Spaß mache, im Code kittens:: zu tippen
  • ThunderKittens verfolgt das Ziel der Einfachheit und bietet vier templatisierte Typen
    • Register-Tiles: 2D-Tensoren im Register-File
    • Register-Vektoren: 1D-Tensoren im Register-File
    • Shared-Tiles: 2D-Tensoren im Shared Memory
    • Shared-Vektoren: 1D-Tensoren im Shared Memory
  • Tiles werden über Höhe, Breite und Layout parametrisiert
  • Register-Vektoren werden über Länge und Layout parametrisiert, Shared-Vektoren nur über die Länge
    • Shared-Vektoren haben in der Regel keine Bank Conflicts
  • Die bereitgestellten Operationen manipulieren Tiles und Vektoren auf Warp-Ebene oder auf Ebene kooperierender Warp-Gruppen
    • Initializer: etwa das Nullsetzen eines Shared-Vektors
    • Unary Ops: etwa exp
    • Binary Ops: etwa mul
    • Zeilen-/Spalten-Operationen: etwa row_sum
  • Da ThunderKittens direkt in CUDA eingebettet ist, scheitert die Abstraktion laut Beschreibung im Unterschied zu Bibliotheken wie Triton „gracefully“
    • Wenn eine Funktion fehlt, lässt sie sich in der gewünschten Weise erweitern

FlashAttention-Beispiel und Leistung

  • Als Beispiel für ThunderKittens wird ein einfacher Forward-FlashAttention-Kernel für die RTX 4090 gezeigt
    • Er behandelt nur headdim=64
    • n muss ein Vielfaches von 256 sein
    • Er umfasst etwa 60 Zeilen CUDA-Code
    • Die Hardware-Auslastung liegt bei 75 %
    • Der Großteil der Komplexität steckt eher im Algorithmus als in Swizzling-Mustern oder Register-Layouts
  • Auch der Forward-Pass von FlashAttention-2 für die H100 wurde mit ThunderKittens geschrieben
    • ThunderKittens kapselt dabei die Komplexität von TMA, WGMMA, Swizzling-Modi und Deskriptoren
    • Der Kernel ist etwa 100 Zeilen lang
    • Auf der H100 ist er etwa 30 % schneller als FlashAttention-2
  • ThunderKittens wird als eine Art „mini-pytorch“ für GPUs beschrieben, das Layouts und Instruktionen kapselt und Primitive bereitstellt
  • Ebenfalls veröffentlicht werden ein Kernel für Based Linear Attention und weitere Kernel für andere Architekturen, die noch folgen sollen
    • Der Based-Linear-Attention-Kernel läuft mit 215 TFLOPs
    • Rechnet man den Recompute des Algorithmus ein, liegt er über 300 TFLOPs
    • Linear Attention ist theoretisch effizienter, war auf realer Hardware historisch aber deutlich weniger effizient
    • Dieses Ergebnis könnte den Einsatzbereich von Anwendungen mit hohem Durchsatz erweitern

Denken in Tiles

  • Dass ThunderKittens gut funktioniert, liegt nach dieser Sichtweise gerade daran, dass es nicht versucht, alles abzudecken
    • CUDA ist deutlich ausdrucksstärker als ThunderKittens
    • ThunderKittens ist ein kleines, einfaches DSL
  • Die zentrale Abstraktion ist das kleine Tile, was gut zur Richtung von AI und Hardware passt
  • ThunderKittens unterstützt keine Dimensionen kleiner als 16
    • Auch die Hardware scheint kein besonderes Interesse an so kleinen Dimensionen zu haben
    • Sinngemäß wird gefragt: Wenn eine Matrixmultiplikation kleiner als 16x16 ist, kann man dann noch sicher sein, dass es überhaupt AI ist?
  • Die aus der CPU-Welt stammende Sicht auf ein 32-Bit-Wort als Register passt nach dieser Auffassung nicht zu AI-Hardware
    • Das 1024-Bit-Vektorregister von CUDA wird als Schritt in die richtige Richtung gesehen
    • Ein Register enthält hier die Daten eines 16x16-Tiles
  • Da AI weiterhin stark von Matrixmultiplikation, Reduction und Reshape geprägt ist, passt die Tile-Abstraktion sowohl zur AI als auch zur Hardware
  • Künftig sollten AI-Ideen so neu geordnet werden, dass sie sich gut auf Hardware abbilden lassen
    • Der recurrent state sollte groß genug sein, um in einen SM zu passen
    • Die Compute-Dichte darf nicht unter dem von der Hardware geforderten Niveau liegen
    • Die auf Hardware gewonnenen Erkenntnisse an AI-Design zurückzukoppeln, wird als wichtiger künftiger Weg gesehen

AMD-Unterstützung geplant

  • Unterstützung für AMD-Hardware in ThunderKittens soll bald erscheinen

Verwandte Beiträge

1 Kommentare

 
GN⁺ 2024-05-13
Hacker-News-Kommentare

  • Die Frage „Wenn die Matrixmultiplikation kleiner als 16x16 ist, bist du dann wirklich sicher, dass es AI ist?“ ist interessant.
    Die Anforderungen an AI-Hardware werden immer klarer. GPUs wurden ursprünglich für einen völlig anderen Zweck entwickelt, wurden aber für AI genutzt, weil ihre Matrixmultiplikations-Hardware gut war, und eine „AI GPU“ könnte einige Funktionen weglassen, die in einer echten GPU vorhanden sind.
    Auch bei der Zahlendarstellung gibt es einen Trend zu kürzeren Formaten wie 16-Bit-Gleitkommazahlen, 8 Bit, 2 Bit und 1 Bit, und irgendwann wird sich ein sinnvoller Punkt einpendeln. Dieser Beitrag zeigt, dass Hardware, die 16x16-Tiles bevorzugt, durchaus plausibel ist. Irgendjemand schreibt so etwas wahrscheinlich schon in VHDL oder wird es bald tun.
    Am Ende scheint es wahrscheinlich, dass einfachere, weniger universelle und billigere Geräte erscheinen werden, die nur „AI“-Berechnungen möglichst ohne unnötigen Hardware-Ballast ausführen

    • GPUs haben sich bereits zu möglichst schlanken AI-Maschinen entwickelt. Zumindest seit der Gründung von Nervana im Jahr 2014 hieß es, GPUs seien alte Technik und für AI ungeeignet, aber offenbar hatte niemand erwartet, dass sie sich so schnell zu AI-Maschinen entwickeln würden
    • Apple geht schon seit einigen Jahren in diese Richtung. Die NPU auf dem Die ist völlig anders als GPU oder CPU[1]
      Nvidia arbeitet vermutlich ebenfalls daran, aber geschäftlich könnte es die bessere Entscheidung sein, an einem Gerät festzuhalten, das Gaming/Entertainment/Krypto/AI bündelt, also an der Form einer Grafikkarte
      [1] https://github.com/hollance/neural-engine/blob/master/docs/a...
    • Die Passage über „NVIDIAs Lügen“ zeigt, wie tief der Wettbewerb geht. Dass ein Dokumentationsfehler völlig zufällig ist, erscheint unwahrscheinlich, und da Diagramme leicht gestohlen oder kopiert werden können, könnte es für Nvidia nützlich gewesen sein, ihn absichtlich stehen zu lassen.
      Das erinnert an die Zeit, als Naveen Raos Nervana einen Nvidia-Maxwell-Treiber baute, der schneller war als Nvidias eigener Treiber. Nicht jeder Doku-Fehler bei einem schnell wachsenden Produkt ist eine Wettbewerbsmaßnahme, aber wenn man bedenkt, dass Forschende lange gebraucht haben, um wgmma per Reverse Engineering zu verstehen, und dazu noch die politische Lage zwischen den USA und China rund um den H100 kommt, wirkt es so, als würde Nvidia alte Methoden nutzen, um seinen Burggraben zu verteidigen.
      Deshalb sollte man sich nicht zu sehr in die Besonderheiten des H100 verbeißen, sondern berücksichtigen, dass zur Frage „Welche Hardware will AI?“ auch die kommerzielle Lage gehört
    • AMD ist mit der Versal-Reihe bereits bei der zweiten Generation angekommen
      https://www.amd.com/en/products/accelerators/alveo/v80.html
      XDNA Architecture
      https://www.amd.com/en/technologies/xdna.html
    • Baut Google solche Geräte nicht auch schon seit fast zehn Jahren?

  • Die Passage „NVIDIAs Lügen. Eine extrem irreführende Darstellung des tatsächlichen 128b-swizzled-wgmma-Layouts. Wegen dieses Diagramms habe ich drei unwiederbringliche Wochen meines Lebens verloren, also stelle ich sie öffentlich bloß“ ist eindrucksvoll.
    Ich frage mich, ob es Leute überrascht, dass ein gewaltiger Teil der Fortschritte bei AI in Ingenieursarbeit wie der Optimierung von Matrixmultiplikationen steckt und dass ein erheblicher Teil dieser Ingenieursarbeit im Reverse Engineering von NVIDIA-Chips besteht

    • Die Architektur selbst macht keinen großen Unterschied. Wenn man ausreichend große Modelle mit ausreichend großen Daten trainiert, kommen unabhängig von der Architektur oft ähnliche Ergebnisse heraus. Deshalb kann man sagen, dass der Großteil des AI-Fortschritts inzwischen darauf beruht, Matrizen sehr schnell multiplizieren zu können

  • Warp-Scheduler, vier Quadranten, Tensor Memory Accelerator, unswizzled wgmma layout…
    Die Grenze zwischen GPU-Terminologie und Star-Trek-artigem Technobabble wird immer unschärfer

    • Beim Lesen wusste ich schon ungefähr, worum es geht, aber „quadrants warpen mit dem Tensorbeschleuniger“ klingt wirklich nach Star Trek.
      Beim Lesen anderer Beiträge hatte ich diesen Gedanken auch schon manchmal. Ich frage mich, wie das auf jemanden wirkt, der hier einen Link zu einem Artikel bekommt und ihn liest. Es fühlt sich wahrscheinlich an, als würde man auf eine Trek-Nerd-Veranstaltung geraten, auf der über Warp-Kerne diskutiert wird
    • Dieser Kommentar ließ mich einen Schritt zurücktreten und die Begriffe mit frischem Blick betrachten, und ich musste lachen, weil es wirklich stimmt

  • Um den Stromverbrauch bei der AI-Inferenz zu senken und die Geschwindigkeit zu erhöhen, scheint ein Wechsel zu analogen Approximationsschaltungen am sinnvollsten zu sein.
    Man braucht nicht unbedingt perfekte Gleitkomma-Multiplikation und -Addition, sondern lediglich eine Vorrichtung, die zwei Eingangsspannungen annimmt und eine Ausgangsspannung liefert, die dem Multiplikationsergebnis hinreichend nahekommt.

    • Ich kenne jemanden, der in diese Richtung arbeitet, und habe gehört, dass die größten Hürden darin bestehen, mit bestehender Chipfertigungstechnologie etwas zu bauen, das analoge Logik ermöglicht, ein Design zu finden, das nicht wie eine Antenne wirkt, und damit umzugehen, dass wegen Fertigungstoleranzen jeder physische Chip anders ausfällt, sodass das auszuführende Modell womöglich pro Chip feinjustiert werden muss.
      Ein großer Vorteil ist, dass man float16 statt über 16 Leitungen durch die Spannung auf nur einer Leitung darstellen würde. Theoretisch könnte sogar eine deutlich höhere Präzision als bei float32 möglich sein. Außerdem ließen sich Werte direkt verbinden, statt sie erst in die arithmetisch-logische Einheit zu laden, sodass Einsparungen bei Die-Fläche und Energie potenziell mehrere Größenordnungen erreichen könnten.
    • Ich denke zwar, dass es noch lange dauert, bis analoge Schaltungen praktisch nützlich werden, aber ein Bereich, in dem man Ungenauigkeit akzeptieren könnte, wären verrauschte digitale Schaltungen.
      Zum Beispiel könnte man akzeptieren, dass sich eines von einer Million Ausgabebits umdreht, wenn sich dadurch das Verhältnis von Leistung zu Stromverbrauch verbessert. Bei float32 wäre das schwierig, weil schon ein einzelner Unendlichkeitswert alles ruinieren kann, aber bei int8 könnte man es vielleicht verkraften, wenn statt 0 gelegentlich 128 herauskommt.
      [1] Ich weiß nicht genau, ob die Matrix-Gleitkommaeinheiten des H100 tatsächlich IEEE 754 einhalten.
    • Einen Schritt weiter gedacht braucht es meiner Meinung nach etwas, das der Funktionsweise eines echten biologischen Gehirns ähnlicher ist und sich dennoch leicht herstellen lässt.
      Biologische neuronale Netze sind nicht annähernd so voll verbunden wie typische künstliche neuronale Netze; die Ein-/Ausgangs-Konnektivität von Neuronen liegt unter 10 und ist damit sehr lokal. In der Biologie gibt es, soweit wir wissen, auch kein Backpropagation, stattdessen aber Feedback und Rekurrenz.
      Es könnte unterstützende Zellen oder Prozesse geben, die für die Funktion des zentralen Nervensystems essenziell sind und die wir noch nicht kennen. Selbst auf hoher Ebene könnte es eine erhebliche Menge „hartkodierter“ Konnektivität geben, und ein Teil davon ist bereits bekannt. Beispielsweise sind die auditiven Neuronen im Ohr verschaltet, und zur Bestimmung der Schallrichtung geschieht dort etwas, das einer Faltung ähnelt. Das ist kein emergentes Phänomen, sondern eine Funktion, die auch ohne Training möglich ist.
      Das Leben hat das über Milliarden Jahre und eine ähnlich große Zahl von Generationen herausgefunden, daher ist es nicht überraschend. Theoretisch wäre das vielleicht auch in Software möglich, aber wenn man die über 1 Billion Neuronen eines Primaten-/Menschengehirns betrachtet, ist das selbst auf heutigen Maschinen mit Tausenden Kernen extrem schwierig. Auch die „Cloud“ würde die dafür nötige Konnektivität und Latenz nicht erfüllen.
      Es wäre cool, wenn sich mit so einem Ansatz erfolgreich etwas auf dem Niveau von Würmern oder Insekten modellieren ließe.
    • Es wirkt fast unmöglich, gleichzeitig genügend Dynamikbereich und Präzision zu erreichen.
    • Ehrlich gesagt wirkt das wie ein Albtraum beim Debugging.

  • Dieser Artikel hat die Freude wieder aufleben lassen, die ich im Kurs CS 149 Parallel Programming empfunden habe.

    • Kayvon und Kunle sind großartig. Ich habe vor zwei Semestern CS149 Parallel Programming belegt, und es war wirklich klasse :)

  • Der Schreibstil dieses Artikels ist wirklich beeindruckend, und ich freue mich darauf, das auf AMD MI300x zu sehen. Wenn jemand möchte, dass ich etwas Zeit auf meiner Hardware investiere, sagt Bescheid.

    • Mich würde interessieren, ob du schon viel AI-Arbeit mit AMD-Produkten gemacht hast. Ich möchte nicht mehr als 2500 Dollar für eine RTX 4090 ausgeben und denke daher über eine RX 7900XTX für Experimente oder den Einstieg nach.
      Mich interessiert, wie gut sie in der Praxis wirklich funktioniert, ob es besser wäre, noch etwas zu sparen und statt der 7900 XT die XTX zu kaufen, und wie stark sich weniger VRAM auf die Nutzbarkeit im Alltag auswirkt.
    • Gute Texte sollten klar und unmissverständlich sein. Beim Sprechen kann man unterbrechen und um clarification bitten, aber ein Text hat nur eine einzige Chance, seine Botschaft zu vermitteln.
      Leser sollten nicht extra auf knowyourmeme.com nachschauen müssen, um zu verstehen, was die Autoren sagen wollen. Ich weiß nicht einmal, was dieser Titel bedeuten soll, und finde, damit verfehlt er sein Ziel deutlich.
    • Wirklich? Das löst bei mir PTSD aus den Wallstreetbets-Zeiten aus.

  • Ich frage mich, wo man anfangen sollte und welcher Roadmap man folgen müsste, um solche Artikel wirklich vollständig zu verstehen.

    • Es gibt eine gute Vorlesung, um GPU-Programmierung zu lernen. Etwa ab Vorlesung 4.0 bekommt man die nötigen Grundlagen: https://youtube.com/playlist?list=PLzn6LN6WhlN06hIOA_ge6Srgd...
      Danach sollte man selbst einmal einen CUDA-Kernel schreiben, der eine Vektor-Matrix-Multiplikation ausführt. Mit pycuda kann man sich auf den Kernel konzentrieren und den Rest in Python schreiben. Man kann ChatGPT sagen, dass man selbst eine Implementierung bauen möchte, die einen Vektor mit 4000 Elementen und eine 4000x12000-Matrix multipliziert, und sich den gesamten Prozess erklären lassen.
      Zum Mieten von GPUs ist Runpod gut; dort gibt es inzwischen alles von günstigen GPUs bis zum H100. Für den Anfang reicht ein günstigeres Modell.
    • Wenn man tiefer einsteigen will, dürfte die Spiral-Playlist zur Matrixmultiplikation hilfreich sein: https://www.youtube.com/playlist?list=PL04PGV4cTuIWT_NXvvZsn...
      Ich habe zwei Monate damit verbracht, mit Spiral Matrixmultiplikations-Kernel zu implementieren und zu optimieren.

  • Das Diagramm im GitHub-README (https://github.com/HazyResearch/ThunderKittens/blob/main/att...) ist viel zu unruhig. Sind solche Wellen-Balkendiagramme überhaupt legal? :P

  • Der Name ThunderKittens ist großartig. Ich würde gern sehen, wie ThunderKittens sich an FlashAttention-Backpropagation wagt, die noch einmal eine Größenordnung schwieriger ist als der Forward-Pass.

  • Ist das nicht die Art von Forschung, die Teams, die heute NPUs bauen, bereits betreiben? Zum Beispiel kann der Groq-Chip seine aktuelle Leistung liefern, weil er eine KI-spezifische Architektur verwendet. Im Consumer-Bereich ist auch Apple Silicon ziemlich leistungsfähig.
    Ich arbeite nicht in diesem Feld, aber es scheint mir Grenzen zu geben, wenn man sich nur auf General-Purpose-Prozessoren stützt, die über vergleichsweise langsame Pfade kommunizieren. Es wirkt sinnvoller, das Design auf Hardware-Ebene neu zu denken und letztlich die Preise im Consumer-Markt zu senken.

    • Bei der Aussage, dass Apple Silicon im Consumer-Bereich ziemlich leistungsfähig sei, bin ich mir nicht so sicher. Im localllama-Subreddit auf Reddit gibt es viele frustrierte Beiträge von CPU-Nutzern, die sich abmühen, überhaupt auf brauchbare Geschwindigkeiten zu kommen.
      Wenn man für ein paar hundert Dollar eine Nvidia-GPU oder für 900 Dollar ein Gaming-Notebook mit einer 4050 und 6 GB VRAM kaufen kann, ist es schwer, CPU-basierte KI als leistungsfähig zu bezeichnen.
      Ich habe es bei der Arbeit auch ohne GPU auf CPU-Basis versucht, aber abgesehen davon, kleine Modelle zu verwenden und zu warten, war es nicht wirklich praktikabel. Am Ende habe ich dann doch einen GPU-Rechner angefordert.
      „Technisch möglich“ und „in der Praxis gut nutzbar“ sind nicht dasselbe. Nvidia ließ sich wirklich gut nutzen, während CPU schmerzhaft und frustrierend war.