Kapitel 5: Entwurf des GPU-Speichersystems

Grafikprozessoren (GPUs) haben sich zu hochparallelen, programmierbaren Beschleunigern entwickelt, die bei einer Vielzahl von Anwendungen eine hohe Leistung und Energieeffizienz erzielen können. Das Speichersystem ist eine entscheidende Komponente moderner GPU-Architekturen, da es die enorme Anzahl gleichzeitiger Threads mit schnellem Zugriff auf Daten versorgen muss. In diesem Kapitel werden wir die Schlüsselelemente des GPU-Speichersystem-Designs untersuchen, einschließlich der in GPUs verwendeten DRAM-Technologien, Speichercontroller und Arbitrierung, gemeinsam genutzter Speicher und Caches sowie Techniken zur effizienten Speichernutzung.

DRAM-Technologien für GPUs

Dynamic Random Access Memory (DRAM) ist die primäre Technologie, die zur Implementierung des Hauptspeichers in modernen Computersystemen, einschließlich GPUs, verwendet wird. DRAM bietet im Vergleich zu anderen Speichertechnologien eine hohe Dichte und relativ geringe Kosten. Allerdings weist DRAM auch eine höhere Zugriffslatenz und geringere Bandbreite im Vergleich zu on-chip Speichern wie Caches und Register-Dateien auf.

GPUs verwenden in der Regel spezialisierte DRAM-Technologien, die auf hohe Bandbreite anstatt auf geringe Latenz optimiert sind. Einige gängige DRAM-Technologien, die in GPUs verwendet werden, sind:

GDDR (Graphics Double Data Rate): GDDR ist eine spezialisierte DRAM-Technologie, die für Grafikkarten und Spielkonsolen entwickelt wurde. Sie bietet eine höhere Bandbreite als Standard-DDR-DRAM durch Verwendung eines breiteren Busses und höherer Taktraten. GDDR5 und GDDR6 sind die neuesten Versionen mit Bandbreiten von bis zu 512 GB/s bzw. 768 GB/s.
HBM (High Bandwidth Memory): HBM ist eine hochleistungsfähige 3D-gestapelte DRAM-Technologie, die eine sehr hohe Bandbreite und niedrigen Stromverbrauch bietet. HBM stapelt mehrere DRAM-Chips übereinander und verbindet sie mit sogenannten Through-Silicon Vias (TSVs), wodurch deutlich höhere Datenübertragungsraten als bei herkömmlichem DRAM möglich sind. HBM2 kann Bandbreiten von bis zu 1 TB/s liefern.

Abbildung 5.1 zeigt den Unterschied zwischen herkömmlichem GDDR-Speicher und 3D-gestapeltem HBM.

   GDDR-Speicher                         HBM-Speicher
  ____________                   ______________________  
 |            |                 |  ___________________  |
 |   DRAM     |                 | |                   | |
 |   Chips    |                 | |    DRAM-Chips     | |
 |            |                 | |___________________| |
 |            |                 |           .          |
 |            |                 |           .          | 
 |            |                 |           .          |
 |____________|                 |  ___________________  |
      |                         | |                   | |
     Leiterplatte                | |  Logik-Chip (GPU) | |
                                | |___________________| |
                                |______________________|

Abbildung 5.1: Vergleich der GDDR- und HBM-Speicherarchitekturen.

Die Wahl der DRAM-Technologie hängt von den spezifischen Anforderungen der GPU ab, wie zum Beispiel dem Strombudget, dem Formfaktor und den Zielanwendungen. Hochleistungs-GPUs für Spiele und professionelle Grafiken verwenden häufig GDDR6 aufgrund ihrer hohen Bandbreite, während HBM2 in Rechenzentren und HPC-GPUs, bei denen Energieeffizienz eine wichtige Rolle spielt, häufiger anzutreffen ist.

Speichercontroller und Arbitrierung

Speichercontroller sind für die Verwaltung des Datenflusses zwischen der GPU und dem DRAM außerhalb des Chips verantwortlich. Sie verarbeiten Speicheranfragen der GPU-Kerne, planen DRAM-Befehle und optimieren Speicherzugriffsmuster, um die Bandbreitenauslastung zu maximieren und die Latenz zu minimieren.

GPU-Speichercontroller verwenden in der Regel ein Mehrkanal-Design, um eine hohe Bandbreite und parallelen Zugriff auf den DRAM bereitzustellen. Jeder Speicherkanal ist mit einem oder mehreren DRAM-Chips verbunden und verfügt über eigene Befehls- und Datenbusse. Der Speichercontroller verteilt Speicheranfragen auf die verfügbaren Kanäle, um Parallelität zu maximieren und Konflikte zwischen den Kanälen zu vermeiden.

Abbildung 5.2 zeigt ein vereinfachtes Diagramm eines GPU-Speichercontrollers mit vier Kanälen.

          GPU-Kerne
              |
        ______|______
       |             |
       |  Speicher   |
       |  Controller |
       |_____________|
         |    |    |    |
        Kanal 0  Kanal 1  Kanal 2  Kanal 3
         |    |    |    |
        DRAM DRAM DRAM DRAM

Abbildung 5.2: GPU-Speichercontroller mit vier Kanälen.

Speicherarbitrierung ist der Prozess der Entscheidung, welche Speicheranfragen zuerst bedient werden sollen, wenn mehrere ausstehende Anfragen vorhanden sind. GPUs verwenden verschiedene Arbitrierungspolicen, um die Leistung und Fairness des Speichersystems zu optimieren:

First-Come, First-Served (FCFS): Die einfachste Arbitrierungspolitik, bei der Anfragen in der Reihenfolge bedient werden, in der sie eintreffen. FCFS ist fair, kann jedoch aufgrund fehlender Anfragenneuordnung zu suboptimaler Leistung führen.
Round-Robin (RR): Anfragen werden zyklisch bedient, wobei allen Anforderern die gleiche Priorität eingeräumt wird. RR gewährleistet Fairness, optimiert jedoch möglicherweise nicht die Lokalität oder Dringlichkeit von Anfragen.
Prioritätsbasiert: Anfragen werden basierend auf verschiedenen Kriterien Prioritäten zugewiesen, wie z. B. der Art der Anfrage (lesen vs. schreiben), der Quelle (Textur vs. L2-Cache) oder dem Alter der Anfrage. Höher priorisierte Anfragen werden zuerst bedient.
Deadline-Aware: Anfragen werden basierend auf ihren Fristen geplant, um eine rechtzeitige Fertigstellung sicherzustellen. Dies ist besonders wichtig für Echtzeit-Grafikanwendungen.
Lokalitätsbewusst: Der Speichercontroller versucht, Anfragen, die auf nahegelegene Speicherorte zugreifen, zusammen zu planen, um Row-Buffer-Treffer zu maximieren und DRAM-Vorladungs- und Aktivierungsüberkopf zu minimieren.

Fortgeschrittene GPU-Speichercontroller setzen oft eine Kombination dieser Arbitrierungspolicen ein, um das beste Gleichgewicht zwischen Leistung, Fairness und Echtzeitanforderungen zu erreichen.

Shared Memory und Caches

GPUs verwenden ein hierarchisches Speichersystem, das sowohl softwaregesteuerte als auch hardwaregesteuerte Caches umfasst, um die Latenz- und Bandbreitenanforderungen des Hauptspeichers zu reduzieren.

Shared Memory

Shared Memory ist ein softwaregesteuerter, on-chip Speicherplatz, der unter den Threads eines Thread Blocks (NVIDIA) oder einer Workgroup (OpenCL) aufgeteilt ist. Es fungiert als nutzerkontrollierter Cache und ermöglicht Programmierern die explizite Verwaltung von Datenbewegungen und -wiederverwendung innerhalb eines Thread Blocks.

Shared Memory wird typischerweise mit schnellen, mehrportigen SRAM-Bänken implementiert, um einen niedrigen Latenz- und hohen Bandbreitenzugriff zu ermöglichen. Jede Bank kann eine Speicheranfrage pro Taktzyklus bedienen, daher muss die Hardware zwischen gleichzeitigen Zugriffen auf dieselbe Bank vermitteln, um Konflikte zu vermeiden.

Abbildung 5.3 veranschaulicht die Organisation des Shared Memory in einem GPU-Kern.

        Thread Block
   ______________________
  |  _________________   |
  | |    Thread 0     |  |
  | |_________________|  |
  |         .            |
  |         .            |
  |         .            |
  |  _________________   |
  | |    Thread N-1   |  |
  | |_________________|  |
  |______________________|
             |
     ________|________
    |                 |
    |  Shared Memory  |
    |  ____________   |
    | | Bank 0     |  |
    | |____________|  |
    | | Bank 1     |  |
    | |____________|  |
    |       .         |
    |       .         |
    |       .         |
    | | Bank M-1   |  |
    | |____________|  |
    |_________________|

Abbildung 5.3: Organisation des Shared Memory in einem GPU-Kern.

Der richtige Einsatz von Shared Memory kann die Leistung von GPU-Kernels erheblich verbessern, indem die Anzahl der Zugriffe auf den langsameren, off-chip DRAM reduziert wird. Es erfordert jedoch eine sorgfältige Programmierung, um effizientes Datenaustausch und Bankkonflikte zu gewährleisten.

Hardwaregesteuerte Caches

Neben dem softwaregesteuerten Shared Memory verwenden GPUs auch hardwaregesteuerte Caches, um automatisch die Datenlokalität auszunutzen und DRAM-Zugriffe zu reduzieren. Die gängigsten Arten von hardwaregesteuerten Caches in GPUs sind:

L1-Daten-Cache: Ein kleiner per-Kern-Cache, der kürzlich abgerufene globale Speicherdaten speichert. Der L1-Cache ist in der Regel privat für jeden GPU-Kern und wird zur Reduzierung der Latenz von globalen Speicherzugriffen verwendet.
Textur-Cache: Ein spezialisierter Cache, der den Zugriff auf schreibgeschützte Texturdaten optimiert. Der Textur-Cache ist für die räumliche 2D-Lokalität optimiert und unterstützt hardwarebeschleunigte Filter- und Interpolationsoperationen.
Konstanten-Cache: Ein kleiner schreibgeschützter Cache, der häufig verwendete Konstantendaten speichert. Der Konstanten-Cache wird an alle Threads in einem Warp übertragen, wodurch er für Daten effizient ist, die von vielen Threads gemeinsam genutzt werden.
L2-Cache: Ein größerer gemeinsam genutzter Cache, der zwischen den GPU-Kernen und dem Hauptspeicher liegt. Der L2-Cache speichert Daten, die aus den L1-Caches ausgelagert werden, und dient zur Reduzierung der Anzahl der DRAM-Zugriffe.

Abbildung 5.4 zeigt eine typische GPU-Speicherhierarchie mit hardwaregesteuerten Caches.

      GPU-Kern 0         GPU-Kern 1         GPU-Kern N-1
   ________________     ________________     ________________
  |                |   |                |   |                |
  |    L1-Daten    |   |    L1-Daten    |   |    L1-Daten    |
  |    Cache       |   |    Cache       |   |    Cache       |
  |________________|   |________________|   |________________|
  |                |   |                |   |                |
  |    Textur      |   |    Textur      |   |    Textur      |
  |    Cache       |   |    Cache       |   |    Cache       |
  |________________|   |________________|   |________________|
  |                |   |                |   |                |
  |    Konstanten  |   |    Konstanten  |   |    Konstanten  |
  |    Cache       |   |    Cache       |   |    Cache       |
  |________________|   |________________|   |________________|
         |                     |                     |
         |_____________________|_____________________|
                               |
                        _______|_______
                       |               |
                       |   L2-Cache    |
                       |_______________|
                               |
Hauptarbeitsspeicher

Abbildung 5.4: GPU-Speicherhierarchie mit hardwaregesteuerten Caches.

Mit hardwaregesteuerten Caches wird die Leistung von GPU-Anwendungen verbessert, indem Datenlokalität automatisch ausgenutzt und die Anzahl der DRAM-Zugriffe reduziert wird. Allerdings können sie auch Cache-Kohärenz und -Konsistenzprobleme verursachen, insbesondere im Kontext von parallelen Programmiermodellen wie CUDA und OpenCL.

Techniken zur effizienten Speicherutilisierung

Eine effiziente Nutzung des GPU-Speichersystems ist entscheidend für hohe Leistung und Energieeffizienz. Einige wichtige Techniken zur Optimierung der Speichernutzung in GPU-Anwendungen sind:

Zusammenfassung: Anordnung von Speicherzugriffen von Threads in einem Warp an benachbarten Speicherpositionen, um sie zu einer einzelnen, breiteren Speichertransaktion zusammenzufassen. Zusammenfassung maximiert die Nutzung der DRAM-Bandbreite und reduziert die Anzahl der Speichertransaktionen.
Optimierung der Datenstruktur: Organisation von Datenstrukturen im Speicher zur Maximierung der räumlichen Lokalität und Minimierung von Cache-Misses. Dies umfasst Techniken wie die Struktur-von-Arrays (SoA)-Anordnung, bei der Datenelemente desselben Typs zusammen gruppiert werden, und die Array-von-Strukturen (AoS)-Anordnung, bei der Datenelemente, die zu derselben Struktur gehören, zusammengehalten werden.
Caching und Vorausladen: Effektive Nutzung von hardwaregesteuerten Caches durch Ausnutzen der zeitlichen und räumlichen Lokalität in den Speicherzugriffsmustern. Dies kann durch Techniken wie Datenzerteilung erreicht werden, bei der Daten in kleinere Chunks aufgeteilt werden, die in den Cache passen, sowie durch Software-Vorausladen, bei dem Daten explizit in den Cache geladen werden, bevor sie benötigt werden.
Speicherzugriffsplanung: Umordnung von Speicherzugriffen zur Maximierung von Zeilenpufferhits und Minimierung von DRAM-Vorkonditions- und -Aktivierungsüberkopf. Dies kann durch Hardware-Mechanismen im Speichercontroller oder durch Software-Techniken wie Optimierung von Zugriffsmustern und Umwandlung von Datenstrukturen erfolgen.
Kompression: Anwendung von Datenkompressionstechniken zur Reduzierung der Größe der zwischen Speicher und GPU-Kernen übertragenen Daten. Dies kann dazu beitragen, Engpässe bei der Bandbreite zu beseitigen und den Energieverbrauch im Zusammenhang mit Datenbewegungen zu reduzieren.
Speichervirtualisierung: Anwendung von Virtual-Memory-Techniken zur Bereitstellung eines vereinheitlichten, zusammenhängenden Adressraums für GPU-Anwendungen. Dies ermöglicht eine flexiblere Speicherverwaltung und ermöglicht Funktionen wie bedarfsorientiertes Paging, das den Speicherbedarf reduzieren und die Systemnutzung verbessern kann.

Abbildung 5.5 zeigt einige dieser Techniken im Kontext eines GPU-Speichersystems.

       GPU-Kerne
          |
    ______|______
   |             |
   | Zusammenfassung |
   |_____________|
          |
    ______|______
   |             |
   | Datenstruktur |
   | Optimierung|
   |_____________|
          |
    ______|______
   |             |
   | Caching und |
   | Vorausladen |
   |_____________|
          |
    ______|______
   |             |
   |   Speicher    |
   |   Zugriffs    |
   | Planung |
   |_____________|
          |
    ______|______
   |             |
   | Kompression |
   |_____________|
          |
    ______|______
   |             |
   |   Speicher    |
   |Virtualisierung|
   |_____________|
          |
        DRAM

Abbildung 5.5: Techniken zur effizienten Speicherutilisierung in einem GPU-Speichersystem.

Zusammenfassung: Anordnung von Speicherzugriffen von Threads in einem Warp an benachbarten Speicherpositionen, um sie zu einer einzelnen, breiteren Speichertransaktion zusammenzufassen. Zusammenfassung maximiert die Nutzung der DRAM-Bandbreite und reduziert die Anzahl der Speichertransaktionen.

Beispiel:
```
// Unzusammengefasstes Zugriffsmuster
int idx = threadIdx.x;
float val = input[idx * stride];
 
// Zusammengefasstes Zugriffsmuster
int idx = threadIdx.x;
float val = input[idx];
```
Optimierung der Datenstruktur: Organisation von Datenstrukturen im Speicher zur Maximierung der räumlichen Lokalität und Minimierung von Cache-Misses. Dies umfasst Techniken wie die Struktur-von-Arrays (SoA)-Anordnung, bei der Datenelemente desselben Typs zusammen gruppiert werden, und die Array-von-Strukturen (AoS)-Anordnung, bei der Datenelemente, die zu derselben Struktur gehören, zusammengehalten werden.

Beispiel:
```
// Array-of-Structures (AoS)-Anordnung
struct Punkt {
    float x;
    float y;
    float z;
};
Punkt points[N];
 
// Structure-of-Arrays (SoA)-Anordnung
struct Punkte {
    float x[N];
    float y[N];
    float z[N];
};
Punkte points;
```

Caching und Vorausladen: Effektive Nutzung von hardwaregesteuerten Caches durch Ausnutzen der zeitlichen und räumlichen Lokalität in den Speicherzugriffsmustern. Dies kann durch Techniken wie Datenzerteilung erreicht werden, bei der Daten in kleinere Chunks aufgeteilt werden, die in den Cache passen, sowie durch Software-Vorausladen, bei dem Daten explizit in den Cache geladen werden, bevor sie benötigt werden.

Beispiel:

// Datenzerteilung
for (int i = 0; i < N; i += TILE_SIZE) {
    for (int j = 0; j < N; j += TILE_SIZE) {
        // Verarbeitung eines Datenkachel, die in den Cache passt
        for (int ii = i; ii < i + TILE_SIZE; ii++) {
            for (int jj = j; jj < j + TILE_SIZE; jj++) {
                // Berechnung von A[ii][jj] durchführen
            }
        }
    }
}

Speicherzugriffsplanung: Umordnung von Speicherzugriffen zur Maximierung von Zeilenpufferhits und Minimierung von DRAM-Vorkonditions- und -Aktivierungsüberkopf. Dies kann durch Hardware-Mechanismen im Speichercontroller oder durch Software-Techniken wie Optimierung von Zugriffsmustern und Umwandlung von Datenstrukturen erfolgen.
Kompression: Anwendung von Datenkompressionstechniken zur Reduzierung der Größe der zwischen Speicher und GPU-Kernen übertragenen Daten. Dies kann dazu beitragen, Engpässe bei der Bandbreite zu beseitigen und den Energieverbrauch im Zusammenhang mit Datenbewegungen zu reduzieren.

Beispiel:
- Delta-Kodierung: Speichern der Unterschiede zwischen aufeinanderfolgenden Werten anstelle der eigentlichen Werte.
- Laufzeichnungskodierung: Ersetzen wiederholter Werte durch eine einzelne Instanz und eine Anzahl.
- Huffman-Kodierung: Zuweisen kürzerer Bit-Sequenzen zu häufiger auftretenden Werten.
Speichervirtualisierung: Anwendung von Virtual-Memory-Techniken zur Bereitstellung eines vereinheitlichten, zusammenhängenden Adressraums für GPU-Anwendungen. Dies ermöglicht eine flexiblere Speicherverwaltung und ermöglicht Funktionen wie bedarfsorientiertes Paging, das den Speicherbedarf reduzieren und die Systemnutzung verbessern kann.

Beispiel:
- Unified Virtual Addressing (UVA) in CUDA: Ermöglicht GPU-Threads den direkten Zugriff auf den CPU-Speicher über einen einzigen Zeiger und vereinfacht die Speicherverwaltung in heterogenen Systemen.

Mehrfach-Chip-Modul-GPUs

Da die Leistungs- und Energieanforderungen von GPUs weiterhin steigen, können herkömmliche Einzelchip-Designs mit der Nachfrage möglicherweise nicht Schritt halten. Mehrfach-Chip-Modul- (MCM-)Designs, bei denen mehrere GPU-Chips in einem einzigen Paket integriert sind, haben sich als vielversprechende Lösung für dieses Problem entwickelt.

MCM-GPU-Designs bieten mehrere Vorteile:

Höhere Speicherbandbreite: Durch Integration mehrerer Speicherstapel oder Chips können MCM-GPUs im Vergleich zu Einzelchip-Designs erheblich höhere Speicherbandbreiten bieten.
Verbesserte Skalierbarkeit: MCM-Designs ermöglichen die Integration von mehr Recheneinheiten und Speichercontrollern, sodass GPUs höhere Leistungsniveaus erreichen können.
Bessere Ausbeute und Kosteneffizienz: Kleinere einzelne Chips in einem MCM-Design können eine bessere Fertigungsausbeute haben und gegenüber großen monolithischen Chips kosteneffizienter sein.

Allerdings führen MCM-GPU-Designs auch neue Herausforderungen mit sich, wie:

Inter-Chip-Kommunikation: Eine effiziente Kommunikation zwischen den verschiedenen Chips in einem MCM-Paket ist entscheidend für die Leistung. Hochbandbreitige, latenzarme Verbindungen sind erforderlich, um den Overhead aus Datenbewegungen zwischen den Chips zu minimieren.
Stromversorgung und thermisches Management: MCM-Designs erfordern sorgfältige Stromversorgungs- und thermische Managementstrategien, um optimale Leistung und Zuverlässigkeit zu gewährleisten.
Softwareunterstützung: MCM-GPUs erfordern möglicherweise Änderungen am Programmiermodell und an den Laufzeitsystemen, um die Vorteile der Mehrchip-Architektur voll auszuschöpfen.

Die Forschung in diesem Bereich untersucht das Design und die Optimierung von MCM-GPUs, einschließlich der Speichersystemarchitektur, dem Verbindungsdesign und dem Ressourcenmanagement.

Ein Beispiel ist das von Arunkumar et al. [2017] vorgeschlagene MCM-GPU-Design, das einen Hochbandbreiten-, latenzarmen Interconnect zur Verbindung mehrerer GPU-Chips verwendet. Die Autoren schlagen auch eine Speichersystemarchitektur vor, die von der erhöhten Bandbreite und Kapazität des MCM-Designs profitiert, um Leistung und Energieeffizienz zu verbessern.

Ein weiteres Beispiel ist die Arbeit von Milic et al. [2018], die ein Ressourcenmanagementkonzept für MCM-GPUs vorschlägt, das darauf abzielt, die Ressourcennutzung zu verbessern und den Overhead der Inter-Chip-Kommunikation zu reduzieren. Das Konzept verwendet eine Kombination aus Hardware- und Softwaretechniken, um die Ressourcennutzung und Kommunikationsmuster der Anwendung zu überwachen und dynamische Ressourcenzuweisungsentscheidungen zu treffen.

Fazit

Das Speichersystem ist eine entscheidende Komponente moderner GPU-Architekturen, und das Design und die Optimierung können einen erheblichen Einfluss auf die Gesamtsystemleistung und -effizienz haben. Da die Anforderungen von parallelen Workloads weiter steigen, erforschen Wissenschaftler eine Vielzahl von Techniken, um die Leistung, Skalierbarkeit und Anpassungsfähigkeit von GPU-Speichersystemen zu verbessern.

Einige wichtige Forschungsrichtungen in diesem Bereich umfassen Speicherzugriffsplanung und Verbindungsdesign, Effektivität von Caches, Priorisierung von Speicheranforderungen und Cache-Umgehung, Ausnutzung von Inter-Warp-Heterogenität, koordinierte Cache-Umgehung, adaptive Cache-Verwaltung, Cache-Priorisierung, Seitenplatzierung im virtuellen Speicher, Datenplatzierung und Mehrfach-Chip-Modul-Designs.Durch die Erforschung dieser und anderer Techniken wollen Forscher GPU-Speichersysteme entwickeln, die den steigenden Anforderungen paralleler Workloads gerecht werden und gleichzeitig hohe Leistung und Energieeffizienz bieten können. Während GPUs weiterentwickelt werden und in Bereichen wie maschinellem Lernen, wissenschaftlicher Berechnung und Datenanalyse neue Anwendungen finden, wird das Design und die Optimierung ihrer Speichersysteme ein wichtiger Bereich der Forschung und Innovation bleiben.

Chapter 4 Gpu Memory System Design Chapter 6 Gpu Performance Metrics and Analysis