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!

joern@dinkla.com

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