Tiny GPU: Eine minimale GPU-Implementierung in Verilog
(github.com/adam-maj)- tiny-gpu ist eine minimale, auf Verilog basierende GPU-Implementierung, mit der man von Grund auf lernen kann, wie eine GPU auf Hardware-Ebene funktioniert, und legt den Fokus eher auf die gemeinsamen Prinzipien von GPGPU und ML-Beschleunigern als auf dedizierte Grafik-Hardware
- Die Implementierung besteht aus weniger als 15 dokumentierten Verilog-Dateien, Architektur- und ISA-Dokumentation, Kernen für Matrixaddition und -multiplikation sowie Unterstützung für Kernel-Simulation und Ausführungstraces
- Die GPU führt jeweils nur einen einzelnen Kernel aus, lädt Programm- und Datenspeicher, setzt
thread_countund startet den Kernel anschließend durch Aktivieren desstart-Signals - Zur Vereinfachung verarbeitet jeder Core jeweils nur einen Block, und jeder Thread besitzt ALU, LSU, PC und Registersatz, wobei angenommen wird, dass alle Threads nach jeder Instruktion wieder auf denselben PC zusammenlaufen
- Die meisten Funktionen moderner GPUs wie mehrstufige Caches, Shared Memory, Memory Coalescing, Pipelining, Warp-Scheduling, Branch Divergence und Barriers wurden weggelassen, um eine lernorientierte Struktur zu priorisieren
Welches Problem tiny-gpu lösen will
- Für CPUs gibt es viele Materialien, mit denen sich von der Architektur bis zu den Steuersignalen lernen lässt, doch die Low-Level-Details moderner GPUs bleiben wegen des wettbewerbsintensiven Markts größtenteils proprietär
- Es gibt viele Materialien zum GPU-Programming, aber kaum Ressourcen, mit denen man lernen kann, wie GPUs auf Hardware-Ebene funktionieren
- Open-Source-GPU-Implementierungen wie Miaow und VeriGPU zielen auf Funktionsvollständigkeit und lauffähiges Verhalten ab und sind daher strukturell komplex
- tiny-gpu entfernt viel von der Komplexität produktionsreifer Grafikkarten und konzentriert sich auf Kernelemente, die modernen Hardware-Beschleunigern gemeinsam sind
- wichtige Komponenten der GPU-Architektur
- wie das SIMD-Programmiermodell in Hardware umgesetzt wird
- wie GPUs mit begrenzter Speicherbandbreite umgehen
Gesamtarchitektur
- tiny-gpu ist so ausgelegt, dass immer nur ein Kernel gleichzeitig ausgeführt wird
- Der Ablauf der Kernel-Ausführung ist wie folgt
- Laden des Kernel-Codes in den globalen Programmspeicher
- Laden der benötigten Daten in den Datenspeicher
- Angabe der Anzahl auszuführender Threads im Device-Control-Register
- Starten der Kernel-Ausführung durch Setzen des
start-Signals auf high
- Die GPU besteht aus folgenden Einheiten
- Device-Control-Register
- Dispatcher
- variable Anzahl an Compute-Cores
- Speichercontroller für Daten- und Programmspeicher
- Cache
Kernel-Ausführung und Thread-Verteilung
- Das Device-Control-Register speichert Metadaten zur Kernel-Ausführung; in tiny-gpu wird darin nur
thread_count, also die Gesamtzahl der auszuführenden Threads, gespeichert - Der Dispatcher verteilt beim Start des Kernels die Threads auf mehrere Compute-Cores
- Er bündelt parallel ausführbare Threads zu Blöcken
- Er weist verfügbare Cores einem Block zur Verarbeitung zu
- Wenn alle Blöcke verarbeitet wurden, signalisiert er den Abschluss der Kernel-Ausführung
- Der vereinfachte Core verarbeitet jeweils nur einen Block
- Jeder Thread besitzt eine eigene ALU, LSU, einen eigenen PC und einen eigenen Registersatz
- Die Verwaltung der Instruktionsausführung der Threads über diese Ressourcen hinweg ist eines der schwierigen Probleme bei GPUs
Speicherstruktur und Controller
- Die GPU ist für die Anbindung an externen globalen Speicher ausgelegt; zur Vereinfachung sind Daten- und Programmspeicher getrennt
- Spezifikation des Datenspeichers
- 8-Bit-Adressierung
- insgesamt 256 Zeilen
- 8-Bit-Daten
- jede Zeile speichert einen Wert kleiner als 256
- Spezifikation des Programmspeichers
- 8-Bit-Adressierung
- insgesamt 256 Zeilen
- 16-Bit-Daten
- gemäß ISA ist jede Instruktion 16 Bit breit
- Der Speichercontroller verfolgt Speicheranfragen der Cores, begrenzt sie entsprechend der tatsächlichen externen Speicherbandbreite und leitet Antworten an die richtigen Ressourcen weiter
- Jeder Speichercontroller besitzt entsprechend der globalen Speicherbandbreite eine feste Anzahl von Kanälen
- Der Cache ist eine Funktion in Arbeit; er speichert aus dem externen Speicher geladene Daten im SRAM des Geräts, damit sie bei späteren Anfragen schneller verfügbar sind und die Speicherbandbreite für neue Daten genutzt werden kann
Interner Aufbau der Cores
- Jeder Core besitzt einen einzelnen Scheduler, der die Thread-Ausführung verwaltet
- Der Scheduler von tiny-gpu führt die Instruktionen eines Blocks vollständig aus, bevor er einen neuen Block holt, und führt die Instruktionen aller Threads in synchronisierter Reihenfolge aus
- In fortgeschritteneren Schedulern können Pipelining und Warp-Scheduling die Ressourcenauslastung erhöhen
- Die wichtigste Einschränkung des Schedulers ist die Latenz beim Laden und Speichern von Daten aus dem globalen Speicher
- die meisten Instruktionen können synchron ausgeführt werden
- Load-Store-Operationen wie
LDRundSTRsind asynchron, daher muss die Instruktionsausführung um lange Wartezeiten herum organisiert werden
- Der Fetcher lädt die Instruktion des aktuellen Program Counters asynchron aus dem Programmspeicher
- Der Decoder dekodiert die geladene Instruktion in Steuersignale für die Thread-Ausführung
- Der Registersatz jedes Threads speichert die Daten der laufenden Berechnung und ermöglicht SIMD-Muster
- schreibgeschützte Register enthalten
%blockIdx,%blockDim,%threadIdx - ein Kernel kann abhängig von der lokalen Thread-ID mit unterschiedlichen Daten ausgeführt werden
- schreibgeschützte Register enthalten
- Die ALU jedes Threads verarbeitet die arithmetischen Instruktionen
ADD,SUB,MUL,DIV CMPgibt aus, ob das Ergebnis der Differenz zweier Register negativ, 0 oder positiv ist, und speichert das Ergebnis imNZP-Register der PC-Einheit- Die LSU jedes Threads greift auf den globalen Datenspeicher zu und verarbeitet
LDR,STRsowie asynchrone Speicherwartezeiten - Der PC jedes Threads bestimmt die als Nächstes auszuführende Instruktion
- standardmäßig wird er pro Instruktion um 1 erhöht
BRnzpspringt zu einer bestimmten Zeile des Programmspeichers, wenn die durch das vorherigeCMPgesetzte Bedingung imNZP-Register erfüllt ist- Schleifen und Bedingungen werden auf diese Weise umgesetzt
- tiny-gpu nimmt zur Vereinfachung an, dass alle Threads nach jeder Instruktion wieder auf denselben PC zusammenlaufen
- In realen GPUs können einzelne Threads zu unterschiedlichen PCs verzweigen; dabei tritt Branch Divergence auf, bei der eine zuvor gemeinsam verarbeitete Thread-Gruppe in mehrere Ausführungsflüsse aufgeteilt wird
ISA
- tiny-gpu implementiert eine ISA mit 11 Instruktionen, um einfache Proof-of-Concept-Kernel wie Matrixaddition und Matrixmultiplikation auszuführen
- Unterstützte Instruktionen
BRnzp: springt zu einer anderen Zeile des Programmspeichers, wenn dieNZP-Bedingung erfüllt istCMP: vergleicht zwei Registerwerte und speichert das Ergebnis imNZP-RegisterADD,SUB,MUL,DIV: grundlegende arithmetische Operationen für Tensor-MathematikLDR: lädt Daten aus dem globalen SpeicherSTR: speichert Daten im globalen SpeicherCONST: lädt einen konstanten Wert in ein RegisterRET: signalisiert das Ende der aktuellen Thread-Ausführung
- Jedes Register wird mit 4 Bit adressiert, insgesamt gibt es also 16 Register
R0bisR12sind 13 frei nutzbare Lese-/Schreibregister- die letzten 3 sind schreibgeschützte Spezialregister für
%blockIdx,%blockDim,%threadIdx, die für SIMD benötigt werden
Ausführungsablauf
- Jeder Core folgt bei der Instruktionsausführung diesem Kontrollfluss
FETCH: nächste Instruktion am aktuellen PC ladenDECODE: Instruktion in Steuersignale dekodierenREQUEST: Daten aus dem globalen Speicher anfordern, fallsLDRoderSTRbenötigt wirdWAIT: falls nötig auf Antwort des globalen Speichers wartenEXECUTE: Berechnungen auf den Daten ausführenUPDATE: Registersatz undNZP-Register aktualisieren
- Dieser Kontrollfluss ist auf Einfachheit und Verständlichkeit ausgelegt
- In einer realen Implementierung könnten einige Schritte zusammengefasst werden, um die Verarbeitungszeit zu optimieren, oder durch Pipelining mehrere Instruktionen koordiniert über die Core-Ressourcen ausgeführt werden
- Jeder Thread führt Berechnungen entlang desselben Ausführungspfads auf Basis der Daten in seinem dedizierten Registersatz aus
- Das ähnelt CPU-Diagrammen, unterscheidet sich aber dadurch, dass
%blockIdx,%blockDim,%threadIdxals schreibgeschützte Register vorhanden sind und so SIMD-Funktionalität ermöglichen
Beispiel-Kernel
- Als Proof of Concept für die ISA wurden Kernel für Matrixaddition und Matrixmultiplikation geschrieben
- Die Testdateien im Repository können diese Kernel vollständig auf der GPU simulieren und den Zustand des Datenspeichers sowie einen vollständigen Ausführungstrace erzeugen
-
Matrixaddition
matadd.asmaddiert zwei 1 x 8-Matrizen- die 8 elementweisen Additionen werden jeweils von einem separaten Thread ausgeführt
- mit den Registern
%blockIdx,%blockDim,%threadIdxwird SIMD-Programmierung demonstriert - mit
LDR- undSTR-Instruktionen wird auch asynchrones Speichermanagement gezeigt
-
Matrixmultiplikation
matmul.asmmultipliziert zwei 2 x 2-Matrizen- das Skalarprodukt der zugehörigen Zeilen und Spalten wird elementweise berechnet
- mit
CMPundBRnzpwird Verzweigung innerhalb eines Threads demonstriert - da alle Verzweigungen wieder zusammenlaufen, funktioniert dies in der aktuellen tiny-gpu-Implementierung
Simulation
- Zum Ausführen der Kernel-Simulation werden iverilog und cocotb benötigt
- Vorbereitung
- Verilog-Compiler und cocotb mit
brew install icarus-verilogundpip3 install cocotbinstallieren - die aktuelle Version von sv2v herunterladen, entpacken und die Binärdatei zu
$PATHhinzufügen - im Repository-Root
mkdir buildausführen
- Verilog-Compiler und cocotb mit
- Die Kernel-Simulation wird mit
make test_mataddundmake test_matmulgestartet - Die Ausführungsergebnisse werden in Logdateien unter
test/logsausgegeben- anfänglicher Zustand des Datenspeichers
- vollständiger Ausführungstrace des Kernels
- finaler Zustand des Datenspeichers
- Am Anfang jeder Logdatei sind die Eingabematrizen zu sehen, am Ende im finalen Datenspeicher die Ergebnis-Matrix
- Der Ausführungstrace enthält für jeden Zyklus den Ausführungszustand aller Threads aller Cores
- aktuelle Instruktion
- PC
- Registerwerte
- Statusinformationen
Absichtlich weggelassene fortgeschrittene GPU-Funktionen
- tiny-gpu lässt zur Vereinfachung die meisten Leistungs- und Funktionsverbesserungen moderner GPUs weg
-
Mehrstufige Caches und Shared Memory
- moderne GPUs verwenden mehrere Cache-Ebenen, um globale Speicherzugriffe zu reduzieren
- tiny-gpu implementiert nur eine einzelne Cache-Ebene zwischen den anfragenden Ressourcen und dem Speichercontroller, die zuletzt verwendete Daten speichert
- mehrstufige Caches cachen häufig genutzte Daten näher am Ort ihrer Verwendung und verkürzen so Ladezeiten
- GPUs verwenden teils auch Shared Memory, damit Threads desselben Blocks gemeinsame Ergebnisse austauschen können
-
Memory Coalescing
- mehrere parallel laufende Threads greifen häufig auf zusammenhängende Adressen zu, etwa auf benachbarte Elemente einer Matrix
- Memory Coalescing analysiert im Queue stehende Speicheranfragen und fasst benachbarte Anfragen zu einer einzigen Transaktion zusammen
- Ziel ist es, den Aufwand für die Adressierung zu verringern und Anfragen gemeinsam zu verarbeiten
-
Pipelining
- im Core von tiny-gpu kann die nächste Instruktion eines Thread-Bündels erst beginnen, wenn die Ausführung der vorherigen abgeschlossen ist
- moderne GPUs streamen mehrere aufeinanderfolgende Instruktionen, während sie für abhängige Instruktionen weiterhin sequentielle Ausführung garantieren
- so wird die Ressourcenauslastung erhöht, damit Core-Ressourcen etwa beim Warten auf asynchrone Speicheranfragen nicht ungenutzt bleiben
-
Warp-Scheduling
- ein Block wird in Warps unterteilt, also in Thread-Gruppen, die gemeinsam ausgeführt werden können
- wenn ein Warp wartet, werden Instruktionen eines anderen Warps ausgeführt, sodass ein einzelner Core mehrere Warps gleichzeitig bearbeiten kann
- das ähnelt Pipelining, betrifft aber Instruktionen unterschiedlicher Threads
-
Branch Divergence
- tiny-gpu nimmt an, dass alle Threads eines einzelnen Batches nach jeder Instruktion denselben PC haben
- in der Praxis können einzelne Threads abhängig von den Daten zu unterschiedlichen Zeilen verzweigen
- Threads mit unterschiedlichen PCs werden in getrennte Ausführungsflüsse aufgeteilt, und auch das spätere Zusammenlaufen muss verwaltet werden
-
Synchronisation und Barriers
- moderne GPUs können Barriers setzen, damit Thread-Gruppen desselben Blocks warten, bis alle einen bestimmten Punkt erreicht haben
- das ist nützlich, wenn Threads gemeinsame Daten austauschen müssen und sichergestellt werden soll, dass die Datenverarbeitung abgeschlossen ist
Nächste Schritte
- Künftige Verbesserungen umfassen
- einfachen Instruktions-Cache hinzufügen
- einen Adapter bauen, um die GPU auf Tiny Tapeout 7 nutzen zu können
- grundlegende Branch Divergence hinzufügen
- grundlegendes Memory Coalescing hinzufügen
- grundlegendes Pipelining hinzufügen
- Kontrollfluss und Registerverwendung zur Verbesserung der Zykluszeit optimieren
- einen einfachen Grafik-Kernel schreiben oder einfache Grafik-Hardware hinzufügen, um Grafikfunktionen zu demonstrieren
- Wer das Repository verbessern möchte, kann per PR beitragen
1 Kommentare
Meinungen auf Hacker News
Eine Ausnahme ist Intel, das viele technische GPU-Dokumente veröffentlicht: https://kiwitree.net/~lina/intel-gfx-docs/prm/
Auch die Handbücher zu i810/815 sind online zu finden, und abgesehen von der seltsamen Lücke im Zeitraum vor 965, in der 855/910/915/945 fehlen, war die Dokumentation ziemlich kontinuierlich.
Dazu gehören sogar Dokumente zur Befehlssatzarchitektur aktueller und früherer Produkte, allerdings wirken sie eher auf Implementierer ausgerichtet als wie High-Level-Erklärungen für interessierte Enthusiasten.
Wünschenswert wäre, dass alle Firmen so vorgingen.
Allerdings sehe ich das eher als SIMD-Koprozessor.
Um es GPU zu nennen, sollte es meiner Meinung nach zumindest irgendeine Form von Display-Ausgabe geben.
Ich weiß, dass der Begriff ziemlich locker geworden ist, seit Nvidia und andere auch reine Server-Varianten von Grafikarchitekturen als GPUs verkaufen, aber der Grafikteil macht bei GPU-Designs auch heute noch einen erheblichen Teil der Komplexität aus.
Eine GPU ohne Ausgabe ist immer noch nützlich.
An meinem Arbeitsplatz gibt es etwa 75 Workstations mit Mittelklasse-Quadro-Karten; die Karten haben nur mini-DisplayPort, während die Firma nur HDMI-Kabel kauft, also sind sie alle an die integrierte Grafik angeschlossen.
Trotzdem beschleunigen diese Karten Software und verarbeiten Grafik, sie geben nur eben kein Bild aus.
Es gibt auch andere Beispiele: https://github.com/jbush001/NyuziProcessor
Welche Stückzahlen wären wohl nötig, damit TSMC oder eine andere Foundry so einen Prozessor wirtschaftlich fertigt?
Ich würde gern mit FPGAs anfangen, aber ehrlich gesagt ist es schon schwer, überhaupt ein Gefühl dafür zu bekommen, wo man anfangen soll, und das ganze Feld wirkt ziemlich einschüchternd.
Mein Endziel wäre eine Beschleunigerkarte für LLMs; auch wenn das völlig willkürlich gewählt ist, dürfte es viele Überschneidungen mit diesem Projekt geben, vermutlich nur mit Unterschieden beim Memory-Offloading, um größere Modelle laden zu können.
Der FPGA-Einstieg muss in mehrere Teiltechnologien zerlegt werden, und man sollte die Erwartungen anpassen.
Von einem Software Engineer erwartet man auch nicht, dass er von den Grundlagen aus einen ganzen Computer baut, eine Befehlssatzarchitektur schreibt, Maschinencode versteht, diesen in Assembly überführt und dann noch eine Programmiersprache entwickelt, um mit Python-Code Anwendungen zu bauen.
Sinnvoll ist es, oben anzufangen und sich im Stack nach unten zu arbeiten.
Wenn man Komplexität abstrahiert und sich darauf konzentriert, Systeme aus vorgefertigter IP zu bauen, ist FPGA-Design ziemlich einfach.
Üblicherweise würde man etwas wie MATLAB empfehlen, weil man dort mit HDL Coder auf einem DevKit mit Reference Design eine erste Anwendung erstellen kann.
Andernfalls entsteht eine enorme Lernlast: digitale Rechnerarchitektur, Verilog, Timing, Transceiver/I/O, Pin-Planung, Quartus/Vivado, Simulation/Verifikation, Embedded Systems und so weiter.
Kurz gesagt: Mit System-Level-Design beginnen, Plug-and-play-IP einbinden, lernen, sie auf oberster Ebene zu verbinden, und dieses Modul dann in ein vorgefertigtes Referenzdesign einsetzen.
Danach kann man nach und nach die Schichten abtragen und die darunterliegende Komplexität freilegen.
Parallel dazu kann man Harris & Harris lesen.
Wenn man die Übungen und das Buch abgeschlossen hat, ist der Zeitpunkt gekommen, ein eigenes Projekt zu starten.
Übrigens gibt es bei HackerMojo auch wöchentliche Treffen; selbst wenn man nicht im Valley ist, kann man per Zoom teilnehmen.
Dort kann man anfangen
Hardware ist tatsächlich der einfache Teil, und zugleich der schwierige Teil, wenn es um die Fertigung geht
always-Block Non-blocking-Zuweisungen und Blocking-Zuweisungen gemischt werden?Es gab eine Website namens opencores, auf der viele Open-Source-HDL-Projekte gesammelt waren
Ich frage mich, ob es heutzutage brauchbare große verteilte HDL-Simulatoren auf HPC-Niveau gibt
Moderne GPUs für Simulationen auf RTL-Ebene zu nutzen, scheint plausibel
Oder ist das nicht dieselbe Website, sondern eine ähnliche andere?
Ist es bei modernen CUDA-Cores und Ähnlichem üblich, Division als echte Instruktion zu haben, oder wird sie normalerweise per Software emuliert?
Eine echte Hardware-Divisionsschaltung braucht enorm viel Fläche, deshalb hätte ich nicht erwartet, dass sie in einer GPU-ALU steckt
In Verilog ist es viel zu leicht, eine Zeile wie
DIV: begin alu_out_reg <= rs / rt; endzu schreiben, aber diese eine Zeile frisst viel SiliziumWenn man nur Verilog simuliert, sieht man das womöglich nicht
Das Projekt bleibt bei der Simulation stehen; um daraus echte Hardware zu machen, wäre deutlich mehr Arbeit nötig
Meiner Meinung nach sollte man so etwas anders nennen
Der Abstand zwischen beiden wird kleiner, und beide Seiten übernehmen Funktionen der jeweils anderen, aber es gibt weiterhin erhebliche Unterschiede
Meiner Ansicht nach hat das mit Amdahls Gesetz zu tun [0]
In diesem Sinne könnte man die CPU als latenzoptimierten Prozessor und die GPU als durchsatzoptimierten Prozessor bezeichnen
Konkreter könnte man [1] die CPU auch als Prozessor für lange, tiefe Datenabhängigkeiten und die GPU als Prozessor für breite, flache Datenabhängigkeiten bezeichnen
[0]: https://en.wikipedia.org/wiki/Amdahl%27s_law
[1]: https://en.wikipedia.org/wiki/Data_dependency
Ein Tensor ist einfach ein n-dimensionales Array
Darauf kann man Software oder Firmware setzen und es wie eine GPU funktionieren lassen
Ich habe versucht, mir aus dem EDK2-Quellcode die Teile zusammenzusetzen, aber es ist unklar, wie viel davon QEMU-spezifisch ist
tiny-gpuannimmt, alle Threads würden nach jeder Instruktion wieder beim selben Program Counter „konvergieren“, ist eine viel zu naive VereinfachungIn echten GPUs können einzelne Threads zu unterschiedlichen PCs verzweigen, und es entsteht Branch Divergence, bei der eine anfangs gemeinsam verarbeitete Thread-Gruppe in getrennte Ausführungen aufgeteilt wird
Es wäre wohl besser gewesen, erst einmal GPU-Programmierung auszuprobieren, bevor man eine GPU in Silizium baut
Außerdem passt auch die Bezeichnung SIMD irgendwie nicht so recht
Das ist dieselbe Person, die früher fremde Schaltungen zusammengesteckt, eine LED zum Blinken gebracht und dann behauptet hat, eine CPU gebaut zu haben
__syncthreads()aufrufen?