CUDA

1. CUDA < Best practices>

2. Compute Unified Device Architecture • Was ist CUDA? • Hardware – Software Architektur • Ermöglicht general-purposecomputing auf einer GPU • 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 • Vorteile • Frei erhältlich • Leicht einzusteigen → ähnliche Syntax wie C

3. Performance - Vergleich • CPU vs. GPU: • Rechenkapazität: 367 GFLOPS vs. 32 GFLOPS • Memory Bandwidth: 86.4 GB/s vs. 8.4 GB/s Floating-Point Operationen pro Sekunde für eine CPU und eine GPU.

4. Struktur – Vergleich (Programmersview) 16 Multiprozessoren (SIMD Einheiten) Control 8 X 32 Bit SP Cache

5. Programmiermodell • 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

6. Anwendungsbeispiel • Matrix – Multiplikation • Thread – Blöcke laden Asub und Bsub • jeder Thread 1 Element • Thread – Block → SubmatrixCsub • 1 Thread → eine Multiplikation

7. Best practices < CUDA>

8. Algorithmus für die GPU optimieren • Daten dürfen nicht von einander abhängig sein • 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 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. Parallelität ausnutzen • Arbeit 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 Multiprozessors • 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

10. “Performance Clip“ • Scenario: • 256 Threads/Block • 3 Blocks pro SIMD - Einheit → 768 Threads (100%) • Wirbenutzen 11 Register pro Thread → 8448 Register • G80 nur 8192 Register vorhanden • Kernel kannnichtgestartetwerden • Lö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: Divergenz vermeiden, wenn Verzweigungsbedingung eine Funktion der Thread-ID ist • 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: • 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  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  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  minimale Latenz (1 Zyklus) • Hunderte mal schneller als Global Memory • Threads desselben Blocks können kommunizieren

13. Global Memory - Coalescing (vereinigt) • Koordiniertes Lesen durch einen Warp (32 Threads) • 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: • 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: das bis zu 10-15 Fache

14. Coalesced Access (Reading floats)

15. Uncoalesced Access (Reading floats)

16. Shared Memory - Bankkonflikte • Shared Memory • 16 KB • organisiert in 16 Banken je 1 KB • Shared Memory • ist genauso schnell wie Register • falls keine Bank Konflikte existieren • 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 (# dergleichzeitigenZugriffe)

17. Shared Memory - Keine Bankkonflikte Zufällige Permutation Broadcast Lineare Adressierung Schrittweite = 3 Wörter Lineare Adressierung Schrittweite = 1 Wort

18. Shared Memory - Bankkonflikte kein Konflikt oder 5 Wege Konflikt Lineare Adressierung Schrittweite = 2 Wörter Lineare Adressierung Schrittweite = 8 Wörter

19. Lösung • Um auf einen 32 Bit Array-Eintrag zuzugreifen wird meistens folgende Syntax benutzt: __shared__ floatshared[32]; float 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__ charshared[32]; char data = shared[BaseIndex + tid]; • Kein Konflikt: char 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 shared[32]; struct type data = shared[BaseIndex + tid]; • Drei separate Speicherzugriffe ohne BK. wenn type so definiert ist: struct type{ float x, y, z; }; • Zwei separate Speicherzugriffe mit BK. wenn type so definiert ist: struct type{ float x, y; }; oder struct type{ float f; char c; };

21. Debugging • CUDA Debuggen → schwer • Auf 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 maximieren • Algorithmus datenparallel strukturieren • Geeignete Konfiguration für maximale Auslastung wählen • Instruktionsnutzung optimieren • Ziel: maximaler Durchsatz 367 GFLOPS • Intrinsic Funktionen statt regulärer Funktionen • Single Precision statt Double Precision • Divergenz bei Flow Control vermeiden • Speicherzugriff 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. → → ? →