Date post: | 05-Apr-2015 |
Category: |
Documents |
Upload: | theudobald-kesting |
View: | 114 times |
Download: | 7 times |
1
CUDA
< Best practices >
2
Compute Unified Device Architecture
Was ist CUDA?Was ist CUDA? Hardware – Software Architektur Ermöglicht general-purpose computing auf einer GPU
Wie funktioniert CUDA?Wie funktioniert CUDA? GPU fungiert als Coprozessor für Haupt-CPU Bearbeitet alle datenparallelen und rechenintensiven Teile
einer Anwendung Diese Programmteile (kernel genannt):
werden in die Instruktionssprache der GPU übersetzt auf der GPU hochgeladen und ausgeführt
VorteileVorteile Frei erhältlich Leicht einzusteigen → ähnliche Syntax wie C
3
Performance - Vergleich
Floating-Point Operationen pro Sekunde für eine CPU und eine GPU.
CPU vs. GPU: CPU vs. GPU:
Rechenkapazität: 367 GFLOPS vs. 32 GFLOPS Memory Bandwidth: 86.4 GB/s vs. 8.4 GB/s
4
Control
Cache
Struktur – Vergleich (Programmers view)
5
Host = CPU Device = GPU
Die CPU leitet Kernel-Aufrufen an
die GPU Jeder Kernel
wird von mehrere Threads bearbeitet
die in einem Gitternetz von Thread-Blöcken organisiert sind
Programmiermodell
6
Anwendungsbeispiel
Matrix – MultiplikationMatrix – Multiplikation Thread – Blöcke laden Asub und Bsub
- jeder Thread 1 Element Thread – Block → Submatrix Csub
- 1 Thread → eine Multiplikation
7
Best practice
s<
CUDA>
8
Algorithmus für die GPU optimieren
Daten dürfen nicht von einander abhängig seinDaten dürfen nicht von einander abhängig sein Maximierung der arithmetischen Intensität (math/bandwidth)Maximierung der arithmetischen Intensität (math/bandwidth)
möglichst viel auf den schon geladenen Daten arbeiten Daten aus dem Speicher holen ist teuer
Auf der GPU arbeitenAuf der GPU arbeiten Strukturen auf Device anlegen, benutzen und löschen Transfer von Daten zwischen Device und Host vermeiden Schwächer parallelisierte Berechnungen können schneller
sein als stark parallelisierte Berechnungen die immer zwischen Device und Host Daten transferieren.
9
Arbeit so einteilen dass die HW voll ausgelastet istArbeit so einteilen dass die HW voll ausgelastet ist Min. so viele Blöcke wie die GPU Multiprozessoren (SIMD -
Einheit) hat verwenden Anzahl der Threads pro Block als Vielfaches der Warp-Size
wählen- besser noch: ein Vielfaches von 64 (laut nVidia)
Je mehr Threads pro Block desto besser Vorsicht: je mehr Threads je weniger Register pro Thread
verfügbar (gleich Beispiel) Konfigurationen mit 100% Auslastung eines Konfigurationen mit 100% Auslastung eines
MultiprozessorsMultiprozessors Ein Multiprozessor kann bis zu 768 Threads bearbeiten
- 2 Blocks x 384 Threads- 3 Blocks x 256 Threads- 4 Blocks x 192 Threads- 6 Blocks x 128 Threads- 8 Blocks x 96 Threads
Parallelität ausnutzen
10
“Performance Clip“
Scenario:Scenario: 256 Threads/Block 3 Blocks pro SIMD - Einheit → 768 Threads (100%) Wir benutzen 11 Register pro Thread → 8448 Register G80 nur 8192 Register vorhanden Kernel kann nicht gestartet werden
LLösung:ösung: Code so ändern dass es nur noch 10 Register pro Threads
braucht → 7680 Register notwendig → Kernel läuft
11
Flow Control – Branch Diverging
Jede Kontroll – Instruktion (if, switch, do, for, while) kann den Instruktionsdurchsatz wesentlich beeinflüssen
Threads nehmen verschiedene Ausführungspfade Diese verschiedene Ausführungspfade werden dann
serialisiert Ein allgemeiner Fall: Ein allgemeiner Fall: Divergenz vermeiden, wenn
Verzweigungsbedingung eine Funktion der Thread-ID ist Beispiel mit Divergenz:Beispiel mit Divergenz:
- if(threadIdx.x > 2){//impl.} else{//impl.} - Erzeugt zwei Ausführungspfade für Threads aus einem Block- Verzweigungsgranularität < Warp Size; - d.h. Threads 0 und 1 folgen einen anderen Pfad als die anderen Threads
aus dem ersten Warp Beispiel ohne Divergenz:Beispiel ohne Divergenz:
- if(threadIdx.x / WARP_SIZE > 2){//impl.} else{//impl}- Erzeugt auch zwei Ausführungspfade für den Threads aus einem Block- Verzweigungsgranularität aber Vielfach der Warp Size; - d.h. alle Threads in einem Warp folgen dem gleichen Pfad.
12
Speichernutzung Host Memory Host Memory maximale Latenz (> 600 Zyklen) maximale Latenz (> 600 Zyklen)
Transfer minimieren- Temporäre Datenstrukturen im Device abspeichern
Gruppentransfer schneller als viele Kleine Datenaustausch über high-performance DMA vom device
Global Memory Global Memory hohe Latenz (400 – 600 Zyklen) hohe Latenz (400 – 600 Zyklen) Zugriff minimieren Typisches Vorgehen:
- Lade Daten aus DRAM in Shared Memory- Synchronisiere Threads zum sicheren Lesen- Bearbeite Daten im Shared Memory- Synchronisiere Threads zum sicheren Zurückschreiben- Schreibe Daten zurück in DRAM
Latenz teilweise versteckt durch Thread-Scheduler Shared Memory Shared Memory minimale Latenz (1 Zyklus) minimale Latenz (1 Zyklus)
Hunderte mal schneller als Global Memory Threads desselben Blocks können kommunizieren
13
Global Memory - Coalescing (vereinigt)
Koordiniertes LesenKoordiniertes Lesen durch einen Warp (32 Threads) Aufeinanderfolgende Speicherbereiche:Aufeinanderfolgende Speicherbereiche:
128 byte – jeder Thread liest ein Wort: int, float 256 byte – jeder Thread liest ein Doppel-Wort: int2, float2 512 byte – jeder Thread liest ein Vierfach-Wort: int4, float4
Einschränkungen:Einschränkungen: Startadresse eines Bereichs muss Vielfaches der Größe
dieses Bereiches sein Der k-te Thread in dem Warp muss auf das k-te Element in
dem gelesenen Speicherbereich zugreifen
nicht alle Threads müssen am Schreib-/Lesevorgang teilnehmen
Gewinn:Gewinn: das bis zu 10-15 Fache
14
Coalesced Access (Reading floats)
15
Uncoalesced Access (Reading floats)
16
Shared Memory - Bankkonflikte
Shared MemoryShared Memory 16 KB organisiert in 16 Banken je 1 KB
Shared Memory Shared Memory ist genauso schnell wie Register falls keine Bank Konflikte existieren
Bank Konflikt: Bank Konflikt: mehrere Threads in der gleichen Halb – Warp greifen auf der
gleiche Bank zu Zugriffe müssen serialisiert werden → Parallelität geht verloren. Kosten = max (# der gleichzeitigen Zugriffe)
17
Shared Memory - Keine Bankkonflikte
Lineare AdressierungSchrittweite = 1 Wort
Zufällige Permutation
Lineare AdressierungSchrittweite = 3 Wörter
Broadcast
18
Shared Memory - Bankkonflikte
Lineare AdressierungSchrittweite = 2 Wörter
Lineare AdressierungSchrittweite = 8 Wörter
kein Konflikt oder 5 Wege Konflikt
19
Lösung
Um auf einen 32 Bit Array-Eintrag zuzugreifen wird meistens folgende Syntax benutzt:__shared__ floatfloat shared[32];floatfloat data = shared[BaseIndex + s * tid]; Wenn die Schrittweite s eine ungerade Zahl ist werden keine
Bankkonflikte entstehen
Zugriff auf Elemente kleiner als 32 Bits: Konflikt:
__shared__ charchar shared[32];charchar data = shared[BaseIndex + tid];
Kein Konflikt:
charchar data = shared[BaseIndex + 4 * tid];
20
Lösung 2
Eine Struktur wird in so viele Memory – Request kompiliert wie es Elemente enthält.
Also wird Folgender Code:__shared__ struct type struct type shared[32];struct type struct type data = shared[BaseIndex + tid]; Drei separate Speicherzugriffe ohne BK. wenn type so definiert
ist: struct typestruct type{ float x, y, z; };
Zwei separate Speicherzugriffe mit BK. wenn type so definiert ist:
struct typestruct type{ float x, y; };
oder
struct typestruct type{ float f; char c; };
21
Debugging CUDA Debuggen → schwerCUDA Debuggen → schwer Auf der CPU debuggen durch emulationAuf der CPU debuggen durch emulation
nvcc –deviceemu (oder linux-Makefile)- man kann alle Funktionen die in C sind aufrufen, sogar in dem Device
Code- der Compiler erkennt Deadlocks
Valgrind benutzen um Speicherzugriffsfehler zu finden- Kann lange dauern- Keine Speicherschutzfunktion auf der GPU
22
Zusammenfassung Parallele Ausführung maximierenParallele Ausführung maximieren
Algorithmus datenparallel strukturieren Geeignete Konfiguration für maximale Auslastung wählen
Instruktionsnutzung optimierenInstruktionsnutzung optimieren Ziel: maximaler Durchsatz 367 GFLOPS Intrinsic Funktionen statt regulärer Funktionen Single Precision statt Double Precision Divergenz bei Flow Control vermeiden
Speicherzugriff optimierenSpeicherzugriff optimieren Ziel: maximale Bandbreite 86,4 GB/s Transfer zw. Host u. Device minimieren Global Memory Zugriff (coalscaled) Shared Memory Zugriff (Bankkonflikte)
Trade-Off zw. Genauigkeit und Geschwindigkeit
23
→ →
→ ?