11 GPU Parallelisierung
# GPU Parallelisierung
Vorteile Actor Model
- Kein Shared Memory
- Keine Data Races
- Distributed Concurrency - Beliebige Verteilung auf mehrere Systeme
- Parallelisierung ist inhärent eingebaut, man braucht keine Synchronisationsprimitiven
Ziel
- GPUs haben viele Cores, aber sie sind langsam und spezifisch (nicht "general-purpose)
- Ziel ist also möglicht hohe Parallelisierung
GPU Prozessarchitektur
- ALU führt Rechenoperationen durch
- Im einer CPU werden relativ wenig Transistoren für die ALU eingesetzt
- Auf einer GPU ist der grösste Teil Transistoren für ALUs, dafür sind diese primitver, z.B. wird nur wenig Caching eingesetzt
- Eine GPU besteht aus mehreren SM (Streaming Multiprocessor), die selbst aus mehreren (z.B. 32 oder 64) Cores bestehen
- die Cores in einer SM sind für sich selbst nicht fähig für Operationen
- Eine Instruktion wird von einer SM ausgeführt, diese wird dann von jedem Core im SM gleichzeitig ausgeführt auf unterschiedlichen Datenstellen
- Normale Operationen werden ausgeführt in sequentiellen Reihenfolge fetch - decode - execute - write back
- GPU-Parallelität eignen sich gut für Blockweise Operationen, wenn viele gleiche Berechnungen durchgeführt werden, z.B. Video Gaming
- Eignet sich hingegen nicht für komplizierte Verzweigungen, z.B. Schach-Computer
NUMA Modell
- Non-Uniform Memory Access
- Kein gemeinsamer Speicher zwischen CPU und GPU
- Unterschiedlicher Instruktionssatz zwischen CPU und CPU
- Code für GPU und CPU sepparat kompilieren
CUDA
- Proprietäre API / Compiler für C/C++ für NVIDIA Karten
-
Video-Computing und CUDA werden parallel in einem Thread-Pool ausgeführt
-
Code für GPU wird speziell ausgezeichnet
- Ein Kernel (Methode) wird von mehreren Threads ausgeführt
- Thread-index wird mit
threadIdx.x
abgerufen (x
ist koordinate)
__global__ void VectorAddKernel(float *A, float *B, float *C) { int i = threadIdx.x; C[i] = A[i] + B[i]; }
Threads
- Threads sind gruppiert in Blöcke, ein Block ist im Hintergrund immer auf einem Streaming Multiprocessor
- Innerhalb eines Blocks können Threads untereinander kommunizieren
- Ein Thread ist ein virtueller Skalarprozessor
- Thread-Blöcke müssen unabhängig sein
- Blöcke werden mit Thread-Pool auf die verschiedenen Streaming-Prozessoren verteilt
- Ein Block ist quasi ein "Task"
- Ein Grid ist eine Gruppe von Blocks, meist ein Programm
- Aufteilung in Blocks und Threads wird vom Programmierer modelliert
- z.B. globalen Index ermittel mit
i = threadIdx.x + blockDim.x * blockIdx.x
- Im Aufruf des Kernels angeben, wie es ausgeführt werden soll.
<<<4, 512>>>>
heisst 4 Blöcke zu je 512 Threads
- z.B. globalen Index ermittel mit
Ausführung
- Speicher muss auf auch der GPU alloziert werden und hin und her von Host- zu Device-Memory kopiert werden
- Diese Schritte müssen manuell im Code gemacht werden
- Resident Blocks & Threads: Maximale Anzahl, die gleichzeitig in der GPU geladen werden kann
CUDA Optimierung
Speichermodell
- Beispiel Matrixmultiplikation: Für jede Zeile in C wird die gleiche Spalte in B verwendet, gleich bei Spalten
- Zugriff auf global Memory ist teuer (600 Zyklen pro Zugriff!)
-
Effizienz-Steigerung mit Caching
-
Im Cuda Speichermodell hat jeder Thread Register zur Verfügung (sehr schnell)
- Threads im gleichen Block können auf ein Shared Memory zugreifen
Tiled Matrix Multiplication
- Problem bei Tiled Matrix Multiplikation: Es kann sein, dass bei grossen Matrizen nicht mal eine ganze Zeile oder Spalte im Cache platz hat!
-
Lösung: In quadratischen Teilen Zwischenresultate berechnen
-
__syncthreads()
dient als Barriere, alle Threads innerhalb eines Blockes werden synchronisiert- Jedes
__syncthreads()
ist eine eigene Barriere! Aufpassen mit Kontrollstrukturen
- Jedes
- Shared-Memory mit
__shared__
deklarieren - Beispiel in Vorlesung: 7 Sek. ohne Optimierung, 0.3 sek. mit Optimierung! -> 150x schneller
Warp
- Ein Block wird intern zu Warps (zu je 32 Threads) zerlegt
- Alle Threads in einem Warp führen gleiche Instruktionen aus (SIMD)
- Ein Block läuft immer auf einem Streaming Prozessor
- Falls ein Warp auf Speicher wartet, führt Streaming Multiprocessor einen anderen Warp aus
- Bei
__syncthreads()
müssen alle Warps im Block aktiv sein!
Divergenz
- Tritt auf, wenn es in einem Warp Abzweigungen (z.B. If.. else) gibt, die nicht garantiert auf einem Warp ausgeführt werden
Memory Coalescing
- Zugriff auf aufeinanderfolgende Daten im Memory werden in einer Transaktion ausgeführt (Memory Burst)
- Ein Memory Burst kann auch von mehreren Threads benutzt werden
- Anwendung: Zugriffe möglichst im Speicher nacheinander ausführen (achten bei Matrix-Multiplikation: Nach Zeilen oder Spalten berechnen)
- Als Offset möglichst den Thread-Index verwenden