Nvidia CUDA Paralelní programování na GPU

Mendelova univerzita v Brně Provozně ekonomická fakulta

Nvidia CUDA Paralelní programování na GPU

2014

O čem to bude... ●

Trocha historie

●

Shadery → Unifikace → GPGPU → CUDA

●

Využití GPGPU

●

GPU a jeho Hardware

●

Nvidia CUDA, OpenCL

●

Jak na GPU programovat –

Podroběnjší pohled na architekturu CUDA

–

Ukázky kódu a porovnání implementací

Jak to začalo ●

●

Nejznámější API pro 3D grafiku: –

Glide (cca 1990) –

–

OpenGL (cca 1992) – Silicon Graphics → Khronos Group

–

Direct3D (cca 1995) – RenderMorphics → Microsoft

3dfx Voodoo (1996)

3dfx → x → Nvidia

Grafická pipeline:

●

●

Vertex Shaders + Pixel (Fragment) Shaders Malé programy se specifickou sadou instrukcí pro výpočet změny geometrie, nebo modifikace barevnosti fragmentů

Shadery ●

Dříve → pevně daný počet Vertex a Pixel shaderů

●

Unfied Shader Model –

●

Co největší sjednocení instrukčních sad pro shadery

Unified Shading Architecture –

Každá shader jednotka je schopná provádět jakýkoliv výpočetní úkol

–

Dynamické škálovaní dle potřeby mezi jednotlivými typy shaderů

–

Dostupné na GPU od Nvidia GeForce 8 a ATI Radeon HD 2000

Shadery – Unified Shading Architecture

GPGPU General-Purpose Computing on Graphics Processing Units ●

●

●

Převedení algoritmů z CPU na GPU → vysoce výkonná vícejádrová zařízení s velkou datovou propustností GPU poskytuje vývojáři paralelní procesory k „obecnému použití“ programovatelné v jazyce C GPU nejčastěji zpracovává vektory (RGB, XYZ) a je pro tento typ práce uzpůsobená –

●

→

Od roku 2002 James Fung (University of Toronto) ve spolupráci s firmou Nvidia publikoval několik článků, které nakonec vedly k vydání Nvidia CUDA v roce 2006 (API které umožňuje pomocí jazyka C vytvářet a spouštět kód na grafikách Geforce 8 a novější) a konkurenčního OpenCL.

Využití GPGPU – PhysX ●

Engine pro simulaci fyziky v reálném čase

●

Dnes jeden z nejpoužívanějších enginů

●

●

Vytvořený firmou Ageia spolu s vlastním HW řešením (PPU – physics processing unit) Ageia zakoupena Nvidií v roce 2008 – PPU výpočty realizovány pomocí GPGPU → Nvidia CUDA

Ukázka simulace pomocí fyzikálního enginu

●

https://www.youtube.com/watch?v=JcgkAMr9r5o

●

https://www.youtube.com/watch?v=143k1fqPukk

Využití GPGPU ●

Vědecké výpočty, simulace, Matlab

●

Klasifikace – neuronové sítě, KNN

●

Zpracování videa a zvuku

●

Bioinformatika, medicínské aplikace

●

Počítačové vidění, zpracování obrazu, OpenCV

●

Kryptografie

●

…

CPU – Architektura

●

Intel Haswell (Core i7)

●

Výrobní proces: 22nm

●

4 Jádra

●

1.4B tranzistorů

GPU – Architektura ●

Nvidia Kepler Geforce 7xx Geforce 8xx

●

Nvidia GK110 →

●

Výrobní proces: 28nm

●

●

15 multiprocesorů (SMX) po 192 CUDA jádrech = 2880 CUDA cores 7.1B transistorů

CPU → GPU

●

●

CPU – minimální odezva s nízkým objemem práce za čas (low latency low throughput processors) GPU – maximální objem práce za čas i s horší odezvou (high latency high throughput processors)

CPU → GPU

●

●

CPU – disponuje velkou cache pamětí a instrukční jednotkou (Control). Zvládne tedy dobře optimalizovat vykonávání instrukcí GPU – jde spíše o hrubou sílu. V rámci jednoho multiprocesoru je používána řídící jednotka pro několik ALU (nazývaných stream procesory) a velmi malá cache paměť, což téměř vylučuje jakékoliv optimalizace vykonávání instrukcí

GPGPU API ●

Nvidia CUDA –

●

OpenCL –

●

Compute Unified Device Architecture

Open Computing Language

Direct Compute

CUDA ●

●

●

●

HEMI Cuda (1970)

●

425 bhp

●

Engine: 426 cu in (6.98 L) Hemi V8

Plymouth Barracuda 440 (1971) 390 bhp Engine: 440cu in (7.2 L) V8

CUDA vs. OpenCL

●

●

●

●

●

●

Pouze pro platformu Nvidia

●

SDK verze 1.0 dostupné od února 2007 (Windows, Linux)

●

První funkční balík na trhu Pokročilejší implementace – množství funkcí, které usnadňují programátorům implementaci (High-level i Low-level API) Přehlednější implementace, debugging jádra a práce s pamětí Proprietární - Freeware

●

●

●

●

Pro většinu platforem Nvidia, AMD, Intel Specifikace verze 1.0 v listopadu 2008 první implementace říjen 2009, IBM Inspirované Nvidia CUDA implementované podobně Spíše Low-level API Složitější debugging jádra a kompexnější práce s pamětí, ale obecnější Součást Khronos Group – Royalty Free Open Standard

CUDA ●

Compute Unified Device Architecture

●

Platforma pro paralelizované výpočty na GPU

●

CUDA SDK – balík potřebných nástrojů pro vývoj

●

Vlastní překladač nvcc pro GPU kód

●

Zpětná binární kompatibilita –

●

Aktuálně dostupná verze 5.0 (GM107, GM108) –

●

Programy pro první G8x GPU by měly fungovat bez problémů i na moderních GPU Nové grafiky od GeForce GTX 750 a řady 8xx

Grafické karty řady Quadro → pro pracovní nasazení, věda, grafika – vysoký výkon, vysoká cena …

CUDA – GFLOP/s

CUDA – GB/s

GPU Archtektura – Multiprocesor (G80) ●

Rozložení na obrázku: Dva spárované multiprocesory složené každý z 8 stream procesorů

●

MP = Multi procesor

●

SP = Stream processor

●

L1/L2 = Cache

●

TF = Texture filtering unit

GPU Architektua – celek (G80)

●

128 Stream procesorů – 16MP x 8SP

CUDA – heterogenní programování ●

●

●

●

Rozlišujeme: –

Host

– zařízení kde běží hlavní program (CPU)

–

Device

– zařízení které spouští vlákna (GPU)

Host a device mají oddělené paměti Je nutné nahrát data ke zpracování z host na device a po výpočtu zkopírovat vypočtená data zpět Kvalifikátory funkcí –

__device__

– Vykonávaná a volaná pouze na device

–

__global__

– Vykonávaná na device volaná z host

–

__host__

– Vykonávaná a volaná pouze na host

CUDA – paměti ●

●

●

●

Registry – extrémně rychlé, přístupné jednotlivými vlákny Sdílená paměť – extrémně rychlá, vysoce paralelní, dostupná pro jednotlivé bloky Globální paměť – dostupná pro všechny, pomalá (400 – 800 cyklů), nevhodná pro přístup z vláken Konstantní paměť – readonly, rychlá odezva a velká propustnost

CUDA – Paměť

CUDA – bloky a vlákna ●

●

●

●

●

●

Vlákna jsou sjednocovány do bloků Blok je spouštěn a na mutiprocesoru Bloky čekají ve frontě ke zpracování na dostupných multiprocesorech Blok – možnost 1D, 2D, 3D indexace Bloky vykonávají stejný kernel Identifikace vlákna: interní proměnná threadIdx

CUDA – škálování ●

Rozložení stejného výpočtu na různých GPU s různým počtem multiprocesorů – automatické

CUDA – mřížka a bloky ●

●

●

●

●

Bloky vláken jsou sjednoceny do mřížky Každá mřížka (Grid) může spouštět rozdílné kernely Mřížka > Blok > Vlákno Identifikace vlákna ve vícerozměrném bloku : threadIdx.x threadIdx.y threadIdx.z Identifikace bloku: interní proměnná blockIdx (blockidx.{x,y})

CUDA – příklad ●

Zpracování obrázku 1024 x 1024 px –

●

●

Stanovíme rozměr Bloku na 16 x 16 → 256 vláken –

Potřebujeme 4096 bloků ((1024x1024) / 256 )

–

dim3 threadsPerBlock(16, 16);

Vytvoříme 2D mřížku (64 x 64 bloků) –

●

Potřebujeme 1 048 576 vláken

dim3 blocksPerGrid(imageWidth / threadsPerBlock.x, imageHeight / threadsPerBlock.y);

Spuštění kernelu –

gpuKernel <<< threadsPerBlock, blocksPerGrid >>();

CUDA – kernel ●

void cpu_soucet() {float vys[], float c1[], float c2[]){ for(int i=0, i<SIZE,i++) vys[i] = c1[i] + c2[i];

} ●

__global__ void gpu_soucet(float *vys, float *c1, float *c2){

int i = threadIdx.x; vys[i] = c1[i] + c2[i]; } ●

Překlad kernelu pomocí NVCC – NVIDIA LLVM-based C/C++ překladače

CUDA – vybrané specifikace

Nvidia CUDA – princip práce ●

●

●

●

Odeslání dat na GPU Spuštění výpočtu Vyčkání na dokončení výpočtu Stažení dat z GPU

CPU vs. GPU – Prahování

CPU vs. GPU – CPU Prahování

CPU vs. GPU – GPU Prahování (CUDA)

CPU vs. GPU – GPU Prahování (OpenCL)

CPU vs. GPU – OpenCV ●

OpenCV –

Open Source Computer Vision

–

Knihovna sdružující funkce pro počítačové vidění a strojové učení

Porovnání implelemtací algoritmu prahování Doba provádění algoritmu prahování v závislosti na rozlišení obrázku a použité implementaci

Děkuji za pozornost...