Tipps & Tricks für den erfolgreichen
Einsatz von GPU-Computing
5. Mai 2014
Jörn Dinkla
Version 1.0
Motivation
Angekündigt „In diesem Vortrag wird ein sequenzielles
Programm schrittweise in ein optimiertes massiv-paralleles GPU-Programm überführt…“
Aber CUDA 6.0
Optimierungsergebnisse Kepler
Didaktisch nicht geeignet Zu viel Fachlogik, zu wenig GPU
Daher mehrere kleine Programme, aber mehr GPU
Klärung
Folien und Sourcecode gibt es
unter ...
http://dinkla.net/parallel2014
Download
Tablet, PC, Großrechner, Cloud,
Spielkonsole, Autos …
GPU-Computing überall
Vom Modell zum Bild ...
Computer-Grafik
Tomas Akenine-Mőller © 2002, Quelle: http://www.realtimerendering.com/
Programmierbare Shader
Massiv Parallel
Kepler GK 110
Z. B. NVIDIA Tesla
AMD FirePro
Eigenschaften ECC Speicher
Größerer Speicher
Bessere Double-Performance
Kein Monitor-Anschluss
Keine Kühlung
Auch spezielle Karten
C ++ 11
C ++
C
Frameworks
Device
WDDM
FrameworkCUDA
OpenCL
C++
AMP
DirectX
AMDTreiber TDD
Thrust C++-
Wrapper
Library
CUDA⊖ Benötigt NVIDIA-Hardware
⊕ Teilweise C++
⊕ Am meisten benutzt
C++ AMP⊖ Benötigt Windows/DirectX, Einschränkungen
⊖ Noch „jung“, keine erprobten Libraries
⊕ C++ 11
OpenCL⊖ Geringer Abstraktionsgrad, C99
⊖ Nicht so viele Libraries wie bei CUDA
⊕ Apple, Intel, AMD
Einschränkungen
Schnelle erste Erfolge
2x – 5x Speedup
Dann wird es schwieriger …
Wissen über die Hardware
notwendig
Speedup
2 3 4 5 6 7 8 9 10 11 …
Schneller, Größer, Besser
Nur wenn notwendig!
Parallelisieren?
1080p Ultra HD / 4K720p576p480p
Nicht das Rad neu erfinden
„best practices“
„think parallel“
Empfehlung:
„Structured Parallel Programming“
Michael McCool et. al.
Intel-lastig, Cilk Plus, TBB
Parallele Patterns
Parallele Patterns
Siehe http://www.parallelbook.com/
map (*2) [1..1024] = [2,4,…,2048]
Datenparallelität: Map
Index 0 1 2 3 4 5 6 7 … 1023
Wert 1 2 3 4 5 6 7 8 … 1024
Ergebnis 2 4 6 8 10 12 14 16 … 2048
Mit Thrust
Mit C++ AMP
Lambda-Ausdruck:
[bindings](params) {
body; }
Bolt
Ähnlicher Abstraktionsgrad
Leider nur für AMD
Projekt nicht aktiv
C++-Bindings
C++ (Nicht C++ 11)
Siehe „Schneller Einstieg in OpenCLmit C++-Bindings“
OpenCL
Abstraktion wichtig für Fortschritt
Java, JVM, Web-Frameworks mit Ruby
Rapid Prototyping
Aber:
Hintergrundwissen bei GPU-Computing notwendig!!!
Was ist hinter dem „Vorhang“ des APIs?
Die beiden Programme haben einen Unterschied! Wer findet ihn?
Schöne neue Welt
De
vic
e
Kernel
sequenced
Kernel
transform
Ho
st
d‘
d d
De
vic
e
Kernel
transform
Ho
st
v
av av
v iota v
Beispiel mit thrust (konstruiert, muss nicht so sein!)
Beispiel mit C++ AMP (konstruiert, muss nicht so sein!)
Host
Host und Device
CPU 1
Core
Global Memory
Bus
GPU / Device
Mem. Ctrl.
Prozessor (SM)
C CC C
C CC C
C CC C
Core
Core Core
Global Memory
8/16
Host und Device
Host
Der „Kopf“ des Ganzen
Management Speicher, Queues, Events
Aufruf von Kerneln
Synchronisation
Device
Traditionell „reine Arbeitspferde“
Basis: BaseBuffer
HostBuffer
Unpinned
Pinned (nicht swapbar)
Lokaler Speicher (NUMA)
DeviceBuffer
Buffer
Host- und Device-Buffer
Versionierung der beiden Seiten
get_host()->incr_version()
Aktualisierung je nach Version
update_host()
update_device()
CUDA 6: Unified Memory
BufferPair
Ho
st
Buffer
De
vic
e
Buffer
1. Übergebe ID als Parameter
2. Hole Daten anhand ID
3. Verarbeitung
4. Speichere Daten anhand ID
SPMD / SIMT
Kernel ruft Funktion op auf
Index-Berechnungen extern
Back to the Basics
Größe / Extension
width, height, depth
index(x,y,z)
in_bounds(x,y,z)
int checked_index(x,y,z)
Extent
0 1 2 3
0 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15
1
2
3
Grund: Kernel in der Praxis
Aus Wilt „The CUDA Handbook“, S. 218
„Index-
Schlacht“
OptimiertNachteil: Speicherorganisation
fest verdrahtet
Vor dem Kernel-Aufruf Wechsel der Abstraktionsebene
Host zu Device
Von C++-Datenstrukturen zu Device-Pointern
Aufruf des Kernels
Kernel-Aufruf
Größe des Thread-Block Für Performance wichtig
Hängt von Hardware ab
Zerlegung der Eingabedaten Thread-Block, Work group, Tile
Grid, NDRange
Beispiel Daten 8x8
Grid 2x2
Block 4x4
NestedExtent
Grid = Data/Block
0 1 2 3 4 5 6 7 0 1 2 3 4 5 6 7
0
1
2
3
4
5
6
7
0 1
2 3
Daten, Pixel, Voxel Grid
kernel<<<g, tb, sm, s>>>(params)
Kernel-Konfiguration
dim3 Grid g
dim3 Thread-Block tb
Größe des Shared-Memory sm
cudaStream_t s
Oft abhängig vom Extent
CudaExecConfig(Extent& e)
CudaExecConfig
Benchmark mit folgenden Funktor
und verschieden Block-Größen
Blockgröße und Performance
Magische 32! Optimum 128?
Warum?
Unterschiedlichen Laufzeiten
Anzahl Threads Zeit [ms]
1 686,449
2 333,055
3 222,076
4 166,660
5 133,771
16 42,532
32 21,878
64 12,491
96 9,269
128 8,306
192 8,937
256 8,617
512 8,476
768 10,123
1024 8,563
Ein Thread-Block wird einem
festen SM zugewiesen
Zerlegung eines Blocks in Warps
32 Threads pro Warp
Kleinste Scheduling-Einheit
Thread-Block-Größe
Vielfaches von 32
Grund: Warp
CUDA Nsight (Visual Studio, Eclipse)
Visual Profiler (Eclipse)
Kommandozeile nvprof
OpenCL Intel Vtune (*)
AMD CodeXL
C++ AMP Visual Studio 2013
Profiling
Computation Bound
Alle Prozessoren 100% ausgelastet
Memory Bound
Bandbreite zum Speicher voll
ausgelastet
Latency Bound
Warten auf die Daten
Arten der Auslastung
Sortiert nach Optimierungspotential
NVVP‘s Guided Analysis
Hier!
Memory Bound
Host
Verschiedene Speicher
CPU 1
Core
L1 / L2
Core
L1 / L2
L3 Cache
Global Memory
Bus
GPU
Global Memory
Constant Texture
Prozessor (SM)
Local / Shared / L1
Registers
C CC C
L2 Cache
8/168-20
192/
288
1600
8000
Mem. Ctrl.
128 Kerne sind Optimum
Einfaches Beispiel,
überraschendes Ergebnis
128 Optimum, aber 192 Kerne
„Viel hilft viel“ stimmt hier nicht
Kernel-Konfiguration
Anzahl der Threads = 32*k
Benchmarks notwendig!
Fazit
Lesezugriffe und Schreibzugriffe Anzahl, Datenmenge, Zugriffsmuster
Operationen Komplexität pro Element / Granularität
O(1)
O(log n)
O(n)
O(n^2)
In der Regel memory bound aber evtl. computation bound
Performance einer Map
Mandelbrot-Kernel
BenchmarkThread-Block Laufzeit [ms]
32, 4 2444,86
64, 4 2453,58
32, 6 2453,65
64, 2 2457,14
32, 5 2465,33
192, 2 2469,59
128, 1 2476,28
64, 3 2479,65
96, 3 2480,85
Occupancy
Unterschiede? Profilen!
Limiter!
Max. 16 Thread-Blöcke „resident“ (*)
Einer aktiv (dunkelgrün)
Ready to go (hellgrün)
Wartend auf Speicher (orange)
Z. B. Occupancy 5/8 = 62.5%
Latency Hiding
W0SMX W1 W2 W3 W4 -- -- --
Berücksichtige (Werte für CC 3.5) Max Threads pro Block 1024
Max Blocks pro SMX 16
Max Warps pro SMX 64
Max Threads pro SMX 2048
Anzahl 32-bit Register pro SMX 64K
Max Anzahl Register pro Thread 255
Max Shared Mem pro SMX 48 KB
Max Shared Mem pro Block 48 KB
Occupancy ist komplex
Occupancy Calculator
128 Threads pro Block
16 Thread-Blöcke pro SMX
2048 Threads pro SMX
192 Threads pro Block
Passen nur 10 TB / SMX
Insg. 1920 Threads pro SMX
Denn 11*192 = 2112 > 2048
128 besser als 192
Computation bound
Guided Analysis
1204,75 GFLOPS
Erreichte Performance
Inhärent sequentiell wg. while
Weitere Optimierung?
ABER:
Laut Werbung: 3977 GFLOPS (GTX 780)
Volle Auslastung einer GPU
Alle Kerne auf allen SMX bearbeiten eine FMA, d.h. warten nicht auf den Speicher
FLOPS = Anzahl der Kerne * Takt * 2
780er: 2304 * 863 Mhz * 2
FMA („Fused multiply-add“) zählt als zwei Operationen
Theoretisches Maximum
Schleifen
Integer-Berechnungen
Speicher-Zugriffe
„Instruktions-Mix“
… schwer zu erreichen
Vergleich 5k*5k, 1000 MaxIter
ProzessorIntel Core i7
3820 3.7 GHz
Intel Core i7
4960HQ
2.60GHz
EVGA 780, 967
MHz Base, 1020
MHz Boost
Laufzeit [ms] 24.249 20.784 39,605
Faktoren Faktoren Faktoren
1 Kern 612,271 524,782 1,000
4 Kerne 153,068 131,196
8 Kerne 76,534 65,598
12 Kerne 51,023 43,732
16 Kerne 38,267 32,799
Vergleich mit CPUs
Kopie zum Device?
Pro Node eine Kopie? Performance?
Hierarchische Datenstrukturen
Möglichkeit
Eigene Speicherverwaltung
Ein Array
new überladen
int index
Statt Pointer
Eine Kopie
Hierarchische Datenstrukturen
cudaMallocManaged()
Suche nach 1…7 in Baum
Parallele Baumsuche
4
2 6
7531
Parallele Baumsuche
Single Instruction!!!
Mask-Bit für jeden Thread im Warp
SIMT - Divergenz
0 1 2 3
int tid = treadIdx.x;
if (tid < 2) {
call_expensive_function()
} else {
call_expensive_function2()
}
Warps Code
Divergenz
Divergenz
While-Schleife statt Rekursion
Keine Divergenz mehr!
Verbesserung
„Dynamic Parallelism“
Ab CC 3.5, CUDA 5.5
Rekursive Kernel
Einschränkungen
Stack-Frames benötigen
Speicherplatz
Rekursionstiefe beschränkt
Rekursive Kernel
Wenig Elemente
„Sparse“ Daten
0 1 2 3 4 5 6 7
0
1
2
3
Divergenz
Komprimieren!
Daten liegen „dicht“
„Keine“ Divergenz mehr!
Minimierung von Kopien!
Pack / Expand
0 1 2 3 4 5 6 7 0 1 0 1
0 0
1 1 1 5
2 2 3 6
3 3 1 2
Maximiere in dieser Reihenfolge
1. Parallelität
2. Speicherdurchsatz
3. Berechnungsdurchsatz
Anleitung: Optimierung
Maximiere Parallelität
Kerne
CPUs
GPUs
SMXs
Kerne
Wegen Kopien
Auf dem Device behalten?
Auslastung gering
Überlappung: Kopie und Kernel
Möglichkeit: Streaming
Kopien müssen asynchron sein
cudaMemcpyAsync(…,stream)
Kernel bekommen Argument
kernel<<<g, b, sm, stream>>>(…)
Synchronisation
Nur, wenn notwendig
cudaStreamSynchronize
Streaming
Aufwand
Umstellen des Codes
Einfügungsreihenfolge
HyperQ ab CC 3.5
Bei CC<3.5 Achtung Performance
Auf richtige Reihenfolge achten
Siehe Literatur
Synchronisation über Events
Streaming
Einsatz mehrerer GPUs
Explizite Auswahl des Devices
CUDA: cudaSetDevice()
OpenCL: clCreateContext()
C++ AMP: accelerator_view als
Argument zu parallel_for_each
Multi-GPU
Kernel und Speicher auf gleichem Device!
=> Partionierung / Strategie
Datenstruktur für Zuordnung
map<partition, device>
Vorsicht bei Bibliotheken
Z. B. Thrust
Bisher nicht komfortabel gelöst
Anforderungen
Multi-GPU
Steuerung durch den Host
Möglichkeiten
1. Ein Thread füllt alle Streams
2. Pro GPU ein Thread für zwei
Streams
3. Pro Stream ein Thread
Performance ist systemabhängig
Multi-GPU
OpenMP
Warte
OpenMP
CPUs nutzen
Auslastung der GPU Parallele Kernel
Streams
Auslastung der SMs Thread-Block
Occupancy
Auslastung der Kerne „Instruction level parallelism“ (ILP)
Maximiere Parallelität
Minimiere Synchronisation
Innerhalb Thread-Block
__syncthreads()
Atomare Zugriffe
atomicAdd()
Speicher
Zwischen Kernel-Aufrufen
Maximiere Parallelität
On-Chip-Speicher
Register, Shared/Local
Cache
Spezial-Speicher
Konstanter Speicher, Texturspeicher
Device-Speicher
Minimiere Kopien
„Coalesced“ Transaktionen
Max. Speicherdurchsatz
Transaktionen sind immer 32, 64
oder 128 Byte
Innerhalb eines Warps
„Zusammengefasst“
„Coalesced“
Speicherdurchsatz
31 63 95 127 159 191 223 255 287 319 351 383
y- statt x-Richtung
Speicherdurchsatz
31 63 95 127 159 191 223 255 287 319 351 383
Minimiere Divergenz
Loop-Unrolling
Berechnen statt Speichern
Arithmetik –fast_math
Präzision vs. Geschwindigkeit
Fusion von Kerneln
Max. Berechnungen
Heuristiken, keine Patentrezepte
Bisher Unterschiede bei jeder
Karten-Generation
Benchmarks sind Pflicht!
Neugierde auch!
Literatur ist inzwischen reichlich
vorhanden
Fazit
Einsteiger
„CUDA by Example“ Sanders, Kandrot
Fortgeschritten
CUDA Programming Cook
CUDA Handbook Wilt
CUDA
Spezialisten GPU Computing Gems
„Jade“ und „Emerald“
Doku „Programming Guide“
„Best Practices Guide“
Web CudaZone bei nvidia.com
Parallel Forall (Blog)
Dokumentation docs.nvidia.com
CUDA
Einsteiger/Fortgeschritten „… in Action“
„… Programming Guide“
„Heterogenous Computing
with OpenCL“
Web Khronos
AMD‘s OpenCL Zone
Apple / Mac OS X
OpenCL
Bisher nur ein Buch
„C++ AMP“
Web
Microsoft
Intel‘s Prototyp/PoC „Shevlin Park“
C++ AMP
Bin beide Tage vor Ort
Sprechen Sie mich an!
Ansonsten
Happy GPU-Computing!
Fragen ?
Freiberuflicher Dipl.-Inform.
Schwerpunkte
Parallele Systeme C++ und Java/JVM
GPU-Computing CUDA, OpenCL, C++ AMP
Weitere Informationen
http://www.dinkla.net
Last but not least