1 / 23

CUDA

CUDA. < Best practices >. Compute Unified Device Architecture. Was ist CUDA? Hardware – Software Architektur Ermöglicht general-purpose computing auf einer GPU Wie funktioniert CUDA? GPU fungiert als Coprozessor für Haupt-CPU

andie
Download Presentation

CUDA

An Image/Link below is provided (as is) to download presentation Download Policy: Content on the Website is provided to you AS IS for your information and personal use and may not be sold / licensed / shared on other websites without getting consent from its author. Content is provided to you AS IS for your information and personal use only. Download presentation by click this link. While downloading, if for some reason you are not able to download a presentation, the publisher may have deleted the file from their server. During download, if you can't get a presentation, the file might be deleted by the publisher.

E N D

Presentation Transcript


  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. → ? →

More Related