Kapitel 7: Design des Streaming-Multiprozessors in der GPU-Entwicklung

Der Streaming-Multiprozessor (SM) ist der Grundbaustein der NVIDIA-GPU-Architekturen. Jeder SM enthält eine Reihe von CUDA-Kernen, die Anweisungen im SIMT-Modus (Single Instruction, Multiple Thread) ausführen. Der SM ist für das Management und die Planung von Warps, den Umgang mit Verzweigungsdivergenz und den schnellen Zugriff auf gemeinsamen Speicher und Caches verantwortlich. In diesem Kapitel werden wir die Mikroarchitektur des SM, einschließlich seiner Pipelines, Warp-Scheduling-Mechanismen, des Register-Datei-Designs und der Organisation des gemeinsamen Speichers und des L1-Caches, untersuchen.

Mikroarchitektur und Pipelines des SM

Der SM ist ein hochgradig paralleler und pipelinierter Prozessor, der darauf ausgelegt ist, Hunderte von Threads gleichzeitig effizient auszuführen. Abbildung 7.1 zeigt ein vereinfachtes Blockdiagramm eines SM in der NVIDIA-Volta-Architektur.

                                 Befehlscache
                                         |
                                         v
                                    Warp-Scheduler
                                         |
                                         v
                               Verteilungseinheit (4 Warps)
                                 |   |   |   |
                                 v   v   v   v
                               CUDA-Kern (FP64/FP32/INT)
                               CUDA-Kern (FP64/FP32/INT)
                               CUDA-Kern (FP64/FP32/INT)
                               ...
                               Tensor-Kern
                               Tensor-Kern
                               ...
                               Lade-/Speichereinheit
                               Lade-/Speichereinheit
                               ...
                               Spezialfunktionseinheit
                                         ^
                                         |
                                Register-Datei (64 KB)
                                         ^
```Geteilter Speicher / L1-Cache (96 KB)

Abbildung 7.1: Vereinfachtes Blockdiagramm eines SM in der NVIDIA Volta-Architektur.

Die Hauptkomponenten des SM umfassen:

Instruktions-Cache: Speichert häufig verwendete Instruktionen, um die Latenz zu reduzieren und den Durchsatz zu verbessern.
Warp-Scheduler: Wählt Warps aus, die zur Ausführung bereit sind, und sendet sie an die verfügbaren Ausführungseinheiten.
Dispatch-Einheit: Ruft Instruktionen für bis zu 4 Warps pro Zyklus ab, decodiert sie und sendet sie an die entsprechenden Ausführungseinheiten.
CUDA-Kerne: Programmierbare Ausführungseinheiten, die eine große Bandbreite an ganzzahligen und Gleitkommaoperationen unterstützen. Jeder SM in Volta enthält 64 CUDA-Kerne.
Tensor-Kerne: Spezialisierte Ausführungseinheiten, die für die Beschleunigung von Deep-Learning- und KI-Workloads entwickelt wurden. Jeder SM in Volta enthält 8 Tensor-Kerne.
Lade-/Speichereinheiten: Verwalten Speicheroperationen, einschließlich Laden und Speichern in den globalen Speicher, den gemeinsamen Speicher und die Caches.
Spezialfunktionseinheiten: Führen transzendentale und andere komplexe mathematische Operationen aus.
Registerdatei: Bietet schnellen Zugriff auf threadprivate Register. Jeder SM in Volta hat eine 64 KB große Registerdatei.
Geteilter Speicher / L1-Cache: Ein konfigurierter Speicherbereich, der als softwaregesteuerter Cache (geteilter Speicher) oder als hardwaregesteuerter L1-Datencache verwendet werden kann.

Die SM-Pipeline ist darauf ausgelegt, den Durchsatz zu maximieren, indem mehrere Warps gleichzeitig ausgeführt und Speicherlatenz verborgen werden können. Abbildung 7.2 zeigt eine vereinfachte Darstellung der SM-Pipeline.

    Instruktions-Fetch
            |
            v
    Instruktions-Decodierung
            |
            v
    Operanden-Sammlung
            |
            v
    Ausführung (CUDA-Kerne, Tensor-Kerne, Lade-/Speichereinheiten, Spezialfunktionseinheiten)
            |
            v
    Rückschreiben

Abbildung 7.2: Vereinfachte SM-Pipeline.

Die Pipelinephasen sind wie folgt:

Instruktions-Fetch: Der Warp-Scheduler wählt einen Warp aus, der zur Ausführung bereit ist.


1. **Instruction Fetch**: Der Prozessor holt die nächste Anweisung für dieses Warp aus dem Befehlscache.

2. **Instruction Decode**: Die abgerufene Anweisung wird dekodiert, um den Operationstyp, die Operanden und die Zielregister zu bestimmen.

3. **Operand Collection**: Die für die Anweisung erforderlichen Operanden werden aus der Registerdatei oder dem gemeinsamen Speicher gesammelt.

4. **Execution**: Die Anweisung wird auf der entsprechenden Ausführungseinheit (CUDA-Kern, Tensor-Kern, Lade-/Speichereinheit oder Spezialfunktionseinheit) ausgeführt.

5. **Writeback**: Das Ergebnis der Ausführung wird in die Registerdatei oder den gemeinsamen Speicher zurückgeschrieben.

Um eine hohe Leistung zu erzielen, setzt die SM verschiedene Techniken ein, um die Ressourcenauslastung zu maximieren und Latenzzeiten zu verbergen:

- **Dual-Issue**: Die SM kann pro Takt zwei unabhängige Anweisungen pro Warp ausgeben, was die Parallelität auf Befehlsebene erhöht.
- **Pipelined Execution Units**: Die Ausführungseinheiten sind pipelined, so dass die SM eine neue Operation in einer Einheit starten kann, bevor die vorherige Operation abgeschlossen ist.
- **Latency Hiding**: Die SM kann zwischen Warps auf Zyklusbasis wechseln, was es ihr ermöglicht, die Latenz von Speicherzugriffen und Operationen mit langer Latenz durch die Ausführung von Anweisungen aus anderen Warps zu verbergen.

Beispiel 7.1 zeigt einen einfachen CUDA-Kernel, der die elementweise Addition zweier Vektoren durchführt.

```cpp
__global__ void vectorAdd(int *a, int *b, int *c, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        c[tid] = a[tid] + b[tid];
    }
}

Beispiel 7.1: CUDA-Kernel für Vektoraddition.

In diesem Beispiel berechnet jeder Thread im Kernel die Summe der entsprechenden Elemente aus den Eingangsvektoren a und b und speichert das Ergebnis im Ausgabevektor c. Die SM führt diesen Kernel aus, indem sie jeden Thread einem CUDA-Kern zuweist und Warps von Threads zur Ausführung auf den verfügbaren Kernen plant. Die Lade-/Speichereinheiten werden verwendet, um die Eingabedaten aus dem globalen Speicher abzurufen und die Ergebnisse zurückzuschreiben.

Warp-Scheduling und Divergenz-Handhabung

EfBitte finden Sie hier die deutsche Übersetzung der Markdown-Datei. Für den Code wurde der Code selbst nicht übersetzt, sondern nur die Kommentare.

Eine effiziente Warp-Planung ist entscheidend, um die Leistung des SM zu maximieren. Der Warp-Scheduler ist dafür verantwortlich, die bereiten Warps auszuwählen und in die verfügbaren Ausführungseinheiten zu übermitteln. Das Hauptziel des Warp-Schedulers ist es, die Ausführungseinheiten ausgelastet zu halten, indem immer Warps zum Ausführen bereitstehen.

Der SM verwendet einen zweistufigen Warp-Scheduling-Mechanismus:

Warp-Scheduling: Der Warp-Scheduler wählt basierend auf einer Scheduling-Richtlinie, wie z.B. Round-Robin oder Oldest-First, die bereiten Warps aus. Die ausgewählten Warps werden dann in die verfügbaren Ausführungseinheiten übermittelt.
Instruktions-Scheduling: Innerhalb jedes Warps plant der SM die Instruktionen basierend auf ihren Abhängigkeiten und der Verfügbarkeit der Ausführungseinheiten. Der SM kann mehrere unabhängige Instruktionen aus demselben Warp in einem Zyklus ausgeben, um die Instruktionsparallelität zu maximieren.

Abbildung 7.3 veranschaulicht den zweistufigen Warp-Scheduling-Mechanismus.

    Warp-Pool
    Warp 1 (Bereit)
    Warp 2 (Wartend)
    Warp 3 (Bereit)
    ...
    Warp N (Bereit)
        |
        v
    Warp-Scheduler
        |
        v
    Verteilungseinheit
        |
        v
    Ausführungseinheiten

Abbildung 7.3: Zweistufiger Warp-Scheduling-Mechanismus.

Eine der Schlüsselherausforderungen beim Warp-Scheduling ist der Umgang mit Verzweigungsdivergenz. Im SIMT-Ausführungsmodell führen alle Threads eines Warps die gleiche Instruktion synchron aus. Wenn ein Warp jedoch auf eine Verzweigungsinstruktion (z.B. eine if-else-Anweisung) trifft, können einige Threads den if-Pfad und andere den else-Pfad nehmen. Dieser Zustand wird als Verzweigungsdivergenz bezeichnet.

Um Verzweigungsdivergenz zu behandeln, verwendet der SM eine Technik namens Prädizierung. Wenn ein Warp auf eine divergente Verzweigung trifft, führt der SM beide Verzweigungspfade nacheinander aus und blendet die Threads aus, die nicht den jeweiligen Pfad nehmen. Die Ergebnisse werden dann mithilfe von Prädikatsregistern kombiniert, um sicherzustellen, dass jeder Thread das korrekte Ergebnis erhält.

Beispiel 7.2 zeigt einen CUDA-Kernel mit einer divergenten Verzweigung.Bitte finden Sie nachfolgend die deutsche Übersetzung der Datei "branch".

__global__ void divergentKernel(int *data, int *result) {
    // Tid ist die globale Thread-ID, berechnet aus Block-ID und Thread-ID im Block
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (data[tid] > 0) {
        // Wenn data[tid] > 0 ist, wird result[tid] = data[tid] * 2 gesetzt
        result[tid] = data[tid] * 2;
    } else {
        // Wenn data[tid] <= 0 ist, wird result[tid] = data[tid] * 3 gesetzt
        result[tid] = data[tid] * 3;
    }
}

Beispiel 7.2: CUDA-Kernel mit divergenter Verzweigung.

In diesem Beispiel kann die Verzweigungsbedingung data[tid] > 0 dazu führen, dass einige Threads in einem Warp den if-Zweig ausführen, während andere den else-Zweig ausführen. Die SM (Streaming Multiprocessor) handhabt diese Divergenz, indem sie beide Pfade sequentiell ausführt und die inaktiven Threads in jedem Pfad maskiert.

Abbildung 7.4 veranschaulicht den Prädikationsprozess für ein Warp mit divergenten Threads.

    Warp (32 Threads)
    Thread 1: data[1] = 5, result[1] = 10
    Thread 2: data[2] = -3, result[2] = -9
    ...
    Thread 32: data[32] = 7, result[32] = 14

    Divergente Verzweigung:
    if (data[tid] > 0) {
        result[tid] = data[tid] * 2;
    } else {
        result[tid] = data[tid] * 3;
    }

    Prädikation:
    Schritt 1: Führe if-Zweig mit Maske aus
        Thread 1: result[1] = 10
        Thread 2: (maskiert)
        ...
        Thread 32: result[32] = 14

    Schritt 2: Führe else-Zweig mit Maske aus
        Thread 1: (maskiert)
        Thread 2: result[2] = -9
        ...
        Thread 32: (maskiert)

    Endgültiges Ergebnis:
    Thread 1: result[1] = 10
    Thread 2: result[2] = -9
    ...
    Thread 32: result[32] = 14

Abbildung 7.4: Prädikationsprozess für ein Warp mit divergenten Threads.

Durch die Verwendung von Prädikation kann die SM Verzweigungsdivergenz handhaben, ohne explizite Verzweigungsanweisungen oder Kontrollflussdivergenzen zu benötigen. Divergente Verzweigungen können sich jedoch immer noch auf die Leistung auswirken, da die SM beide Pfade sequentiell ausführen muss, was die effektive Parallelität reduziert.

Registerdatei und Operanden-Collector

Die Registerdatei ist eine kritische Komponente der SM, die einen schnellen Zugriff auf thread-private Register ermöglicht. Jede SM verfügt über eine große Registerdatei, um die vielen aktiven Threads zu unterstützen und einen effizienten Kontextwechsel zwischen Warps zu ermöglichen.Hier ist die deutsche Übersetzung der Markdown-Datei, wobei der Codeblock nicht übersetzt wurde, sondern nur die Kommentare:

In der NVIDIA Volta-Architektur hat jeder SM eine 64 KB große Registerdatei, die in 32 Bänke à 2 KB organisiert ist. Die Registerdatei ist so konzipiert, dass sie eine hohe Bandbreite und eine geringe Latenz bietet, um die große Anzahl von parallelen Threads zu unterstützen.

Um Bankkonflikte zu minimieren und die Leistung zu verbessern, verwendet der SM eine Technik namens Operand-Sammlung. Operand-Sammler sind spezialisierte Einheiten, die Operanden aus den Registerdatei-Bänken sammeln und sie an die Ausführungseinheiten übermitteln. Durch den Einsatz von Operand-Sammlern kann der SM die Auswirkungen von Bankkonflikten reduzieren und die Auslastung der Ausführungseinheiten verbessern.

Abbildung 7.5 zeigt ein vereinfachtes Diagramm der Registerdatei und der Operand-Sammler in einem SM.

    Registerdatei (64 KB)
    Bank 1 (2 KB)
    Bank 2 (2 KB)
    ...
    Bank 32 (2 KB)
        |
        v
    Operand-Sammler
        |
        v
    Ausführungseinheiten

Abbildung 7.5: Registerdatei und Operand-Sammler in einem SM.

Die Operand-Sammler funktionieren, indem sie Operanden aus mehreren Anweisungen und mehreren Warps sammeln, so dass der SM Anweisungen aus verschiedenen Warps in einem einzigen Zyklus an die Ausführungseinheiten senden kann. Dies hilft, die Latenz von Registerdatei-Zugriffen zu verbergen und den Gesamtdurchsatz des SM zu verbessern.

Beispiel 7.3 zeigt einen CUDA-Kernel, der ein Skalarprodukt zweier Vektoren berechnet.

__global__ void dotProduct(float *a, float *b, float *result, int n) {
    // Code-Block
}

In diesem Beispiel berechnet jeder Thread eine Teillösung des Skalarprodukts mit seinem zugewiesenenHier ist die deutsche Übersetzung der Datei:

Elemente aus den Eingangsvektoren. Die Teilergebnisse werden im gemeinsamen Speicherarray partialSum gespeichert. Nachdem alle Threads ihre Teilergebnisse berechnet haben, wird eine parallele Reduktion durchgeführt, um die Teilergebnisse zu summieren und das endgültige Punktprodukt-Ergebnis zu erhalten.

Der Operandensammler spielt in diesem Beispiel eine entscheidende Rolle, indem er die Operanden für den gemeinsamen Speicherzugriff und die arithmetischen Operationen effizient sammelt. Er hilft, Bankenkonflikte zu vermeiden und die Auslastung der Ausführungseinheiten zu verbessern.

Zusammenfassung

Der Streaming-Multiprozessor ist die zentrale Recheneinheit in modernen GPU-Architekturen. Sein Design konzentriert sich darauf, den Durchsatz zu maximieren und die Speicherlatenzen durch eine Kombination aus feingranularer Multithreading-Ausführung, SIMT-Ausführung und effizienter Operandensammlung zu verbergen.

Zu den Schlüsselkomponenten des SM gehören der Warp-Scheduler, der Warps für die Ausführung auswählt; der SIMT-Stack, der Zweigaufteilungen und -konvergenzen handhabt; das Register-File und die Operandensammler, die einen schnellen Zugriff auf threadprivate Register ermöglichen; sowie der gemeinsame Speicher und der L1-Cache, die einen niedrig latenten Datenaustausch und -wiederverwendung ermöglichen.

Da sich GPU-Architekturen weiterentwickeln, wird die Forschung in Bereichen wie Zweigaufteilungshandhabung, Warp-Scheduling und Register-File-Design entscheidend für die Verbesserung der Leistung und Effizienz zukünftiger GPUs sein. Neuartige Techniken wie dynamische Warp-Bildung, Thread-Block-Kompaktierung und Operanden-Wiederverwendungscache haben das Potenzial, die Fähigkeiten des SM deutlich zu verbessern und neue Leistungsniveaus in parallelen Computing-Workloads zu ermöglichen.

Chapter 6 Gpu Performance Metrics and Analysis Chapter 8 Interconnect and on Chip Networks