Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooooo oooooo oooooooooooo oooooooooo oo Uvod, základy CUDA Jiří Filipovič jaro 2013 Jih Filipovič Uvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ♦OOOOOOOO OOOOOOO OOOOOOOOO OOOOOO OOOOOOOOOOOO OOOOOOOOOO OO Motivace - Moorův zákon loorův zákon Počet tranzistorů na jednom čipu se přibližně každých 18 měsíců zdvojnásobí. i -00.0 Jih Filipovič Uvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ♦OOOOOOOO OOOOOOO OOOOOOOOO OOOOOO OOOOOOOOOOOO OOOOOOOOOO OO Motivace - Moorův zákon loorův zákon Počet tranzistorů na jednom čipu se přibližně každých 18 měsíců zdvojnásobí. Adekvátní růst výkonu je zajištěn: • dříve zvyšováním frekvence, instrukčním paralelismem, out-of-order spouštěním instrukcí, vyrovnávacími pamětmi atd. • dnes vektorovými instrukcemi, zmnožováním jader i -00.0 Jih Filipovič Uvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr OÍOOOOOOO ooooooo ooooooooo oooooo oooooooooooo oooooooooo oo Motivace - změna paradigmatu Důsledky Moorova zákona: • dříve: rychlost zpracování programového vlákna procesorem se každých 18 měsíců zdvojnásobí • změny ovlivňují především návrh kompilátoru, aplikační programátor se jimi nemusí zabývat • dnes: rychlost zpracování dostatečného počtu programových vláken se každých 18 měsíců zdvojnásobí • pro využití výkonu dnešních procesorů je zapotřebí paralelizovat algoritmy • paralelizace vyžaduje nalezení souběžnosti v řešeném problému, což je (stále) úkol pro programátora, nikoliv kompilátor JiFÍ Filipovič Úvod, základy CUDA 3/56 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr oo«oooooo ooooooo ooooooooo oooooo oooooooooooo oooooooooo oo Motivace - druhy paralelismu • úlohový paralelismus • problém je dekomponován na úlohy, které mohou být prováděny souběžně • úlohy jsou zpravidla komplexnější, mohou provádět různou činnost o vhodný pro menší počet výkonných jader • zpravidla častější (a složitější) synchronizace • datový paralelismus • souběžnost na úrovni datových struktur • zpravidla prováděna stejná operace nad mnoha prvky datové struktury • jemnější paralelismus umožňuje konstrukci jednodušších procesorů JiFÍ Filipovič Úvod, základy CUDA 4/56 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooo»ooooo ooooooo ooooooooo oooooo oooooooooooo oooooooooo oo Motivace - druhy paralelismu • z pohledu programátora • rozdílné paradigma znamená rozdílný pohled na návrh algoritmů • některé problémy jsou spíše datově paralelní, některé úlohově • z pohledu vývojáře hardware • procesory pro datově paralelní úlohy mohou být jednodušší • při stejném počtu tranzistorů lze dosáhnout vyššího aritmetického výkonu • jednodušší vzory přístupu do paměti umožňují konstrukci HW s vysokou paměťovou propustností i -00.0 Uvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr oooo«oooo ooooooo ooooooooo oooooo oooooooooooo oooooooooo oo Motivace - grafické výpočty datově paralelní • provádíme stejné výpočty pro různé vertexy, pixely, . předdefinované funkce programovatelné funkce • specifické grafické efekty • GPU se stávají stále více programovatelnými • díky tomu lze zpracovávat i jiné, než grafické úlohy i -00.0 Uvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooo»ooo ooooooo ooooooooo oooooo oooooooooooo oooooooooo oo Motivace - výkon Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr oooooo«oo ooooooo ooooooooo oooooo oooooooooooo oooooooooo oo Motivace - výkon Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooo»o ooooooo ooooooooo oooooo oooooooooooo oooooooooo oo Motivace - shrnutí • GPU jsou výkonné • řádový nárůst výkodu již stoji za studium nového programovacího modelu • pro plné využití moderních GPU i CPU je třeba programovat paralelně • paralelní architektura GPU přestává být řádově náročnější • GPU jsou široce rozšířené • jsou levné • spousta uživatelů má na stole superpočítač i -00.0 Uvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr OOOOOOOO* ooooooo ooooooooo oooooo oooooooooooo oooooooooo oo Motivace - uplatnění Využití GPU pro obecné výpočty je dynamicky se rozvíjející oblast s širokou škálou aplikací Jih Filipovič Uvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr OOOOOOOO* OOOOOOO OOOOOOOOO OOOOOO OOOOOOOOOOOO OOOOOOOOOO oo Motivace - uplatnění Využití GPU pro obecné výpočty je dynamicky se rozvíjející oblast s širokou škálou aplikací • vysoce náročné vědecké výpočty • výpočetní chemie • fyzikální simulace • zpracování obrazů • a mnohé další... i -00.0 Jih Filipovič Uvod. základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr OOOOOOOO* ooooooo ooooooooo oooooo oooooooooooo oooooooooo oo Motivace - uplatnění Využití GPU pro obecné výpočty je dynamicky se rozvíjející oblast s širokou škálou aplikací • vysoce náročné vědecké výpočty • výpočetní chemie • fyzikální simulace • zpracování obrazů • a mnohé další... • výpočetně náročné aplikace pro domácí uživatele • kódování a dekódování multimediálních dat • herní fyzika • úprava obrázků, 3D rendering • atd... JiFÍ Filipovič Úvod, základy CUDA 10/56 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr OOOOOOOOO «000000 ooooooooo oooooo oooooooooooo oooooooooo oo Architektura GPU CPU vs. GPU • jednotky jader vs. desítky multiprocesorů • out of order vs. in order • MIMD, SIMD pro krátké vektory vs. SIMT pro dlouhé vektory • velká cache vs. malá cache, často pouze pro čtení GPU používá více tranzistorů pro výpočetní jednotky než pro cache a řízení běhu => vyšší výkon, méně univerzální i -00.0 Uvod. základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo o»ooooo ooooooooo oooooo oooooooooooo oooooooooo oo Architektura GPU Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo oo«oooo ooooooooo oooooo oooooooooooo oooooooooo oo Architektura GPU V rámci systému: • koprocesor s dedikovanou pamětí • asynchronní běh instrukcí • připojen k systému přes PCI-E i -00.0 Uvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr OOOOOOOOO OOO0OOO ooooooooo oooooo oooooooooooo oooooooooo oo Procesor G80 G80 • první CUDA procesor • obsahuje 16 m u Iti procesorů • m u Iti procesor • 8 skalárních procesorů • 2 jednotky pro speciální funkce • až 768 threadů • HW přepínání a plánování threadů • thready organizovány po 32 do warpů • SIMT • nativní synchronizace v rámci multiprocesoru JiFÍ Filipovič Úvod, základy CUDA 14/56 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo oooocoo ooooooooo oooooo oooooooooooo oooooooooo oo Paměťový model G80 Paměťový model • 8192 registrů sdílených mezi všemi thready multiprocesoru • 16 KB sdílené paměti • lokální v rámci multiprocesoru • stejně rychlá jako registry (za dodržení určitých podmínek) • paměť konstant • cacheovaná, pouze pro čtení • paměť pro textury • cacheovaná, 2D prostorová lokalita, pouze pro čtení • globální paměť • pro čtení i zápis, necacheovaná • přenosy mezi systémovou a grafickou pamětí přes PCI-E JiFÍ Filipovič Úvod, základy CUDA 15/56 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr OOOOOOOOO OOOOO^O ooooooooo oooooo oooooooooooo oooooooooo oo Procesor G80 Device Multiprocessor N Multiprocessor 2 Multiprocessor 1 JiFÍ Filipovič Úvod, základy CUDA 16/56 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr OOOOOOOOO 000000« ooooooooo oooooo oooooooooooo oooooooooo oo Další vývoj Procesory odvozené od G80 • double-precision výpočty • relaxovány pravidla pro efektivní přístup ke globální paměti • navýšeny on-chip zdroje (více registrů, více threadů na MP) • lepší možnosti synchronizace (atomické operace, hlasování warpů) Fermi • vyšší paralelizace na úrovni multiprocessoru (více jader, dva warp schedulery, více DP výkonu) • konfigurovatelná LI a sdílená L2 cache • plochý adresní prostor • lepší přesnost v plovoucí řádové čárce • paralelní běh kernelů • širší možnosti synchronizace • další změny plynoucí z odlišné architektury^ JiFÍ Filipovič Úvod, základy CUDA 17/56 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód OOOOOOOOO OOOOOOO »00000000 oooooo oooooooooooo CUDA podrobně oooooooooo Závěr 00 Srovnání teoretické rychlosti GPU a CPU Teoretická maxima • GPU má cca lOx rychlejší aritmetiku • GPU má cca 5x vyšší propustnost paměti • zajímavé pro mnohé problémy (budu čekat na výsledky simulace měsíc nebo rok? pojede mi hra na 3 nebo 30fps?) Některé publikace ukazují lOOx i lOOOx zrychlení • v pořádku, je-li interpretováno jako zrychlení oproti produkčnímu SW (ten nemusí být perfektně optimalizovaný) • interpretováno jako srovnání CPU a GPU zpravidla nesmysl Srovnáváme-li přínos GPU oproti CPU, musíme uvažovat efektivní implementaci pro obě platformy. JiFÍ Filipovič Úvod, základy CUDA 18/56 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód ooooooooo ooooooo o«ooooooo oooooo oooooooooooo CUDA podrobně oooooooooo Závěr 00 Srovnání teoretické rychlosti GPU a CPU V praxi máme však často sériový CPU kód • běh v jednom vlákně znamená až 16x zpomalení (16-jádrové CPU) • absence vektorizace znamená až 4x zpomalení (32-bit operace u SSE instrukcí), 8x u AVX instrukcí Oproti sériové implementaci tedy můžeme kód paralelizací a vektorizací zrychlit • 32x pro čtyřjádrové CPU s AVX nebo osmijádrové s SSE GPU akcelerací pak • cca 300 x Vektorizace a paralelizace pro CPU je však programátorskou náročností srovnatelná s GPU akcelerací. JiFÍ Filipovič Úvod, základy CUDA 19/56 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo oo»oooooo oooooo oooooooooooo oooooooooo oo Teoretické vs. dosažitelné zrychlení Výkonový odstup GPU může být vyšší • jednotky pro speciální funkce, operace na texturách • SIMT pružnější než SIMD • neduhy SMP (omezení škálování propustnosti paměti, „vytloukání řádků cache") Stejně jako nižší • nedostatek paralelismu • příliš vysoký overhead • nevhodný algoritmus pro GPU architekturu Dále se podíváme, jak rozlišit, jestli je nebo naopak není váš algoritmus vhodný pro GPU. JiFÍ Filipovič Úvod, základy CUDA 20/56 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooo«ooooo oooooo oooooooooooo oooooooooo oo Paralelizace Sčítání vektorů • jednoduché datově-paralelní vyjádření • žádná synchronizace • potřebujeme velké vektory Game of Life • co chceme paralelizovat? Game of Life - zjištění nového stavu hry • pro větší herní plochy dostatek paralelismu • jednoduchá synchronizace Game of Life - zjištění stavu buňky po n krocích • inherentně sekvenční? (Game of Life je P-complete, P = A/C) • neznáme paralelní algoritmus JiFÍ Filipovič Úvod, základy CUDA 21/56 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo oooo«oooo oooooo oooooooooooo oooooooooo oo Paralelizace Redukce • na první pohled může vypadat sekvenčně • ve skutečnosti realizovatelná v logn krocích • často je třeba nedržet se sekvenční verze a zamyslet se nad paralelizací problému (ne sekvenčního algoritmu) JiFÍ Filipovič Úvod, základy CUDA 22/56 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooo«ooo oooooo oooooooooooo oooooooooo oo Paralelizace Problém nalezení povodňové mapy • máme výškovou mapu terénu, přítok vody, a chceme zjistit, jaká oblast se zatopí • sekvenčnost dána rozléváním vody • je snadné najít úlohově-paralelní algoritmus, datově-paralelní už tak ne • periodická aktualizace stavu každého bodu mapy o aktualizace omezená jen na hranice vodní plochy (šetří procesory) • rozlévání vody zametači přímkou (vhodnější pro GPU, jednodušší synchronizace) • hledání souvislých oblastí a jejich spojování (odstraňuje sekvenčnost rozlévání) • vždy práce navíc oproti sekvenční/úlohově-paralelní verzi • úkol PV197 na podzim 2010, výkon odevzdaných implementací se lišil o 4 řády (!) JiFÍ Filipovič Úvod, základy CUDA 23/56 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo oooooo«oo oooooo oooooooooooo oooooooooo oo Divergence kódu Divergence kódu • serializace, divergují-li thready uvnitř warpu • nalezení nedivergujícího algoritmu může být snadné • redukce • ale také může prakticky znemožnit akceleraci některých jinak dobře paralelizovatelných algoritmů • mnoho nezávislých stavových automatů • nutnost zamyslet se nad výrazně odlišným algoritmem pro daný problém JiFÍ Filipovič Úvod, základy CUDA 24/56 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooo«o oooooo oooooooooooo oooooooooo oo Divergence přístupu do paměti Divergence přístupu do paměti • není-li do paměti přistupováno po souvislých blocích v rámci warpu, snižuje se její propustnost • často velmi těžko překonatelný problém • průchod obecného grafu • může vyžadovat využití odlišných datových struktur • práce s řídkými maticemi • u rigidnějších struktur si lze často pomoci on-chip pamětí • transpozice matic Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr OOOOOOOOO OOOOOOO OOOOOOOO* oooooo oooooooooooo oooooooooo oo Latence GPU GPU je dnes často propojena se zbytkem systému přes PCI-E • kopírování vstupů/výstupů je relativně pomalé • akcelerovaný algoritmus musí provádět dostatečné množství aritmetiky na přenášená data • násobení matic je vhodné (0(n3) operací na 0(n2) dat) • sčítání vhodné není (0(n2) operací na 0(n2) dat), může být však součástí většího problému i -00.0 Uvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr OOOOOOOOO OOOOOOO OOOOOOOOO «00000 oooooooooooo oooooooooo oo CUDA CUDA (Compute Unified Device Architecture) • architektura pro paralelní výpočty vyvinutá firmou NVIDIA • poskytuje nový programovací model, který umožňuje efektivní implementaci obecných výpočtů na GPU • je možné použít ji s více programovacími jazyky OpenCL Fortran C++ CUDA Architecture m -00,0 Jih Filipovič Uvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr OOOOOOOOO OOOOOOO OOOOOOOOO o*oooo oooooooooooo oooooooooo oo C for CUDA C for CUDA přináší rozšíření jazyka C pro paralelní výpočty • explicitně oddělen host (CPU) a device (GPU) kód • hierarchie vláken • hierarchie pamětí • synchronizační mechanismy • API JiFÍ Filipovič Úvod, základy CUDA 28/56 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooooo oo»ooo oooooooooooo oooooooooo oo Hierarchie vláken Hierarchie vláken • vlákna jsou organizována do bloků • bloky tvoří mřížku • problém je dekomponován na podproblémy, které mohou být prováděny nezávisle paralelně (bloky) » jednotlivé podproblémy jsou rozděleny do malých částí, které mohou být prováděny kooperativně paralelně (thready) • dobře škál uje JiFÍ Filipovič Úvod, základy CUDA 29/56 Motivace Architektura GPU Algoritmy a GPU ooooooooo ooooooo ooooooooo CUDA Demonstrační kód ooo*oo oooooooooooo CUDA podrobně oooooooooo Závěr 00 Hierarchie vláken Grid Block (O, O) Block (1,0) Block (2,0) Block (1,1) Thread (0, 0) Thread (1, 0) 1 Thread (2, 0) Thread (3, 0) Thread (0,1) Thread (1,1) Thread (2,1) Thread (3,1) Thread (0, 2) Thread (1, 2) i Thread (2, 2) Thread (3, 2) JiFÍ Filipovič Úvod, základy CUDA 30/56 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooooo ooooco oooooooooooo oooooooooo oo Hierarchie pamětí Více druhů pamětí • rozdílná viditelnost • rozdílný čas života • rozdílné rychlosti a chování • přináší dobrou škálovatelnost i -00.0 Uvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooooo ooooo* oooooooooooo oooooooooo oo Hierarchie pamětí JiFÍ Filipovič Uvod, základy CUDA 32/56 Algoritmy a GPU OOOOOOOOO Demonstrační kód •ooooooooooo CUDA podrobně oooooooooo Příklad - součet vektorů Chceme sečíst vektory a a b a výsledek uložit do vektoru c. í -00.0 Uvod, základy CUDA Algoritmy a GPU OOOOOOOOO Demonstrační kód •ooooooooooo CUDA podrobně oooooooooo Příklad - součet vektorů Chceme sečíst vektory a a b a výsledek uložit do vektoru c. Je třeba najít v problému paralelismus. i -00.0 Uvod, základy CUDA Algoritmy a GPU OOOOOOOOO Demonstrační kód •ooooooooooo CUDA podrobně oooooooooo Příklad - součet vektorů Chceme sečíst vektory a a b a výsledek uložit do vektoru c. Je třeba najít v problému paralelismus. Sériový součet vektoru: for (int i — 0; i < N; i++) c ľ i 1 = a í i 1 + b f 1 1 : í -00.0 Uvod. základy CUDA Algoritmy a GPU OOOOOOOOO Demonstrační kód •ooooooooooo CUDA podrobně oooooooooo Příklad - součet vektorů Chceme sečíst vektory a a b a výsledek uložit do vektoru c. Je třeba najít v problému paralelismus. Sériový součet vektoru: for (int i — 0; i < N; i++) c[i] = a[i] + b[i]; Jednotlivé iterace cyklu jsou na sobě nezávislé paralelizovat, škáluje s velikostí vektoru. lze je i -00.0 Uvod. základy CUDA Algoritmy a GPU OOOOOOOOO Demonstrační kód •ooooooooooo CUDA podrobně oooooooooo Příklad - součet vektorů Chceme sečíst vektory a a b a výsledek uložit do vektoru c. Je třeba najít v problému paralelismus. Sériový součet vektorů: for (int i — 0; i < N; i++) c[i] = a[i] + b[i]; Jednotlivé iterace cyklu jsou na sobě nezávislé - lze je paralelizovat, škáluje s velikostí vektoru, i-tý thread sečte i-té složky vektorů: c[i] = a[i] + b[i]; Jak zjistíme, kolikátý jsme thread? i -00.0 Uvod. základy CUDA Motivace Architektura GPU Algoritmy a GPU ooooooooo ooooooo ooooooooo CUDA Demonstrační kód oooooo o»oooooooooo CUDA podrobně oooooooooo Závěr 00 Hierarchie vláken Grid Block (O, O) Block (1,0) Block (2,0) Block (1,1) Thread (0, 0) Thread (1, 0) 1 Thread (2, 0) Thread (3, 0) Thread (0,1) Thread (1,1) Thread (2,1) Thread (3,1) Thread (0, 2) Thread (1, 2) i Thread (2, 2) Thread (3, 2) JiFÍ Filipovič Úvod, základy CUDA 34/56 Algoritmy a GPU OOOOOOOOO Demonstrační kód oo«ooooooooo CUDA podrobně oooooooooo Identifikace vlákna a bloku C for CUDA obsahuje zabudované proměnné: • threadldx.jx, y, z} udává pozici threadu v rámci bloku • blockDim. (x, y, z} udává velikost bloku • blockldx.jx, y, z} udává pozici bloku v rámci mřížky (zje vždy 1) • gridDim.jx, y, z} udává velikost mřížky (zje vždy 1) í -00.0 Uvod, základy CUDA Algoritmy a GPU OOOOOOOOO Demonstrační kód ooo»oooooooo CUDA podrobně oooooooooo Příklad - součet vektorů Vypočítáme tedy globální pozici threadu (mřížka i bloky jsou jednorozměrné): i -00.0 Uvod, základy CUDA Algoritmy a GPU OOOOOOOOO Demonstrační kód ooo»oooooooo CUDA podrobně oooooooooo Příklad - součet vektorů Vypočítáme tedy globální pozici threadu (mřížka i bloky jsou jednorozměrné): int i = blockldx.x*blockDim.x + threadldx.x; Uvod, základy CUDA Algoritmy a GPU OOOOOOOOO Demonstrační kód ooo»oooooooo CUDA podrobně oooooooooo Příklad - součet vektorů Vypočítáme tedy globální pozici threadu (mřížka i bloky jsou jednorozměrné): int blockldx.x*blockDim. threadldx.x; Celá funkce pro paralelní součet vektorů: _global__ void addvec(float *a, float *b, float *c){ int i = blockldx.x*blockDim.x + threadldx.x; cíil = aíil + bíil; i -00.0 Uvod, základy CUDA Algoritmy a GPU OOOOOOOOO Příklad - součet vektorů Demonstrační kód ooo»oooooooo CUDA podrobně oooooooooo Vypočítáme tedy globální pozici threadu (mřížka i bloky jsou jednorozměrné): int blockldx.x*blockDim.x threadldx.x; Celá funkce pro paralelní součet vektorů: __global__ void addvec(float *a, float *b , float * c ) { int i = blockldx.x*blockDim.x + threadldx.x; c[i] = a[i] + b[i]; } Funkce definuje tzv. kernel, při volání určíme, kolik threadů a v jakém uspořádání bude spuštěno. Uvod. základy CUDA Algoritmy a GPU OOOOOOOOO Demonstrační kód oooo«ooooooo CUDA podrobně oooooooooo Kvantifikátory typů funkcí Syntaxe C je rozšířena o kvantifikátory, určující, kde se bude kód provádět a odkud půjde volat: • __device__ funkce je spouštěna na device (GPU), lze volat jen z device kódu • __global__ funkce je spouštěna na device, lze volat jen z host (CPU) kódu • __host__ funkce je spouštěna na host, lze ji volat jen z host • kvantifikátory __host__ a __device__ lze kombinovat, funkce je pak kompilována pro obojí Uvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooooo oooooo ooooo«oooooo oooooooooo oo Ke kompletnímu výpočtu je třeba: JiFÍ Filipovič Úvod, základy CUDA 38/56 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooooo oooooo ooooo«oooooo oooooooooo oo Ke kompletnímu výpočtu je třeba: • alokovat paměť pro vektory, naplnit je daty i -00.0 Uvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooooo oooooo ooooo«oooooo oooooooooo oo Ke kompletnímu výpočtu je třeba: • alokovat paměť pro vektory, naplnit je daty • alokovat paměť na GPU Uvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooooo oooooo ooooo«oooooo oooooooooo oo Ke kompletnímu výpočtu je třeba: • alokovat paměť pro vektory, naplnit je daty • alokovat paměť na GPU • zkopírovat vektory a a b na GPU Úvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooooo oooooo ooooo«oooooo oooooooooo oo Ke kompletnímu výpočtu je třeba: • alokovat paměť pro vektory, naplnit je daty • alokovat paměť na GPU • zkopírovat vektory a a b na GPU • spočítat vektorový součet na GPU i -00.0 Uvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooooo oooooo ooooo«oooooo oooooooooo oo Ke kompletnímu výpočtu je třeba: • alokovat paměť pro vektory, naplnit je daty • alokovat paměť na GPU • zkopírovat vektory a a b na GPU • spočítat vektorový součet na GPU • uložit výsledek z GPU paměti do c i -00.0 Uvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooooo oooooo ooooo«oooooo oooooooooo oo Ke kompletnímu výpočtu je třeba: • alokovat paměť pro vektory, naplnit je daty • alokovat paměť na GPU • zkopírovat vektory a a b na GPU • spočítat vektorový součet na GPU • uložit výsledek z GPU paměti do c • použít výsledek v c :-) i -00.0 Uvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU ooooooooo ooooooo ooooooooo CUDA Demonstrační kód oooooo oooooo«ooooo CUDA podrobně oooooooooo Závěr 00 Příklad - součet vektorů CPU kód naplní a a b, vypíše c: (f include Sdefine N 64 int main(){ float a[N] , b[N] , c [N] ; for (int i = 0; i < N; i++) = = i; // zde bude kód provádějící výpočet na GPU for (int i = 0; i < N; i++) printf("%f , " , c[i ] ); return 0; í -00.0 Jih Filipovič Uvod, základy CUDA Motivace Architektura GPU ooooooooo ooooooo Algoritmy a GPU OOOOOOOOO Demonstrační kód ooooooo«oooo CUDA podrobně oooooooooo Správa GPU paměti Paměť je třeba dynamicky alokovat. cudaMalloc(void** devPtr, size_t count); Alokuje paměť velikosti count, nastaví na ni ukazatel devPtr. Uvolnění paměti: cudaFree(void* devPtr); Kopírování paměti: cudaMemcpy(void* dst, const void* src, size_t count, enum cudaMemcpyKind kind); Kopíruje count byte z src do dst, kind určuje, o jaký směr kopírování se jedná (např. cudaMemcpyHostToDevice, nebo cudaMemcpyDevice ToHost). Uvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr OOOOOOOOO OOOOOOO OOOOOOOOO OOOOOO OOOOOOOOÍOOO oooooooooo oo Příklad - součet vektorů Alokujeme paměť a přeneseme data: float *d_a, *d_b, *d_c; cudaMalloc((void**)&d_a, N*sizeof(*d_a)); cud aM alloc (( vo id **)&id_b , N*sizeof(*d_b)); cudaMalloc((void**)&d_c, N*sizeof(*d_c)); cudaMemcpy(d_a, a, N*sizeof(*d_a), cudaMemcpyHostToDevice); cudaMemcpy(d_b, b, N*sizeof(*d_b), CudaMemcpyHostToDevice); // zde bude spuštěn kernel cudaMemcpy(c, d_c, N*sizeof(*c), cudaMemcpyDeviceToHost); cudaFree(d_a); cudaFree(d_b); cudaFree(d_c); JiFÍ Filipovič Úvod, základy CUDA 41/56 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooooo oooooo ooooooooo«oo oooooooooo oo Příklad - součet vektorů Spuštění kernelu: • kernel voláme jako funkci, mezi její jméno a argumenty vkládáme do trojitých špičatých závorek velikost mřížky a bloku • potřebujeme znát velikost bloků a jejich počet • použijeme ID blok i mřížku, blok bude pevné velikosti • velikost mřížky vypočteme tak, aby byl vyřešen celý problém násobení vektorů Pro vektory velikosti dělitelné 32: Sdefine BLOCK 32 addvec«(d_a , d_b , d_c ) ; Jak řešit problém pro obecnou velikost vektoru? JiFÍ Filipovič Úvod, základy CUDA 42/56 Algoritmy a GPU OOOOOOOOO Demonstrační kód oooooooooo»o CUDA podrobně oooooooooo Příklad - součet vektorů Upravíme kód kernelu: __global__ void addvec(float *a, float *b, float *c, int n){ int i = blockldx.x*blockDim.x + threadldx.x; if (i < n) c[i] = a[i] + b[i]; } A zavoláme kernel s dostatečným počtem vláken: addvec«(d_a , d_b , d_c , N); i -00.0 Uvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr OOOOOOOOO OOOOOOO OOOOOOOOO OOOOOO OOOOOOOOOOO* oooooooooo oo Příklad - spuštění Nyní už zbývá jen kompilace :-). nvcc -I/usr/local/cuda/include -L/usr/local/cuda/lib -lcudart \ -o vecadd vecadd.cu Kde s CUDA pracovat? • vlastní stroj: stáhněte a nainstalujte CUDA toolkit a SDK z developer.nvidia.com • windowsí stanice v učebnách (titan) • ke vzdálené práci s high-end GPU: na přání JiFÍ Filipovič Úvod, základy CUDA 44/56 Algoritmy a GPU OOOOOOOOO Demonstrační kód oooooooooooo CUDA podrobně •OOOOOOOOO Paměti lokální v rámci threadu Registry • nejrychlejší paměť, přímo využitelná v instrukcích • lokální proměnné v kernelu i proměnné nutné pro mezivýsledky jsou automaticky v registrech • pokud je dostatek registrů • pokud dokáže kompilátor určit statickou indexaci polí • mají životnost threadu (warpu) i -00.0 Uvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr OOOOOOOOO OOOOOOO OOOOOOOOO OOOOOO OOOOOOOOOOOO »000000000 oo Paměti lokální v rámci threadu Registry • nejrychlejší paměť, přímo využitelná v instrukcích • lokální proměnné v kernelu i proměnné nutné pro mezivýsledky jsou automaticky v registrech • pokud je dostatek registrů • pokud dokáže kompilátor určit statickou indexaci polí • mají životnost threadu (warpu) Lokální paměť • co se nevejde do registrů, jde do lokální paměti • ta je fyzicky uložena v DRAM, je tudíž pomalá a má dlouhou latenci • má životnost threadu (warpu) JiFÍ Filipovič Úvod, základy CUDA 45/56 Algoritmy a GPU OOOOOOOOO Demonstrační kód oooooooooooo CUDA podrobně o«oooooooo Paměť lokální v rámci bloku Sdílená paměť • u c.c. 1.x rychlá jako registry • nedojde-li ke konfliktům paměťových bank • instrukce umí využít jen jeden operand ve sdílené paměti (jinak je třeba explicitní load/store) • v C for CUDA deklarujeme pomocí __shared__ • proměnná ve sdílené paměti může mít dynamickou velikost (určenou při startu), pokud je deklarována jako extern bez udání velikosti pole • má životnost bloku i -00.0 Uvod. základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooooo oooooo oooooooooooo oo»ooooooo oo Paměť lokální pro GPU Globální paměť • řádově nižší přenosová rychlost než u sdílené paměti • latence ve stovkách GPU cyklů • pro dosažení optimálního výkonu je třeba paměť adresovat zarovnaně • má životnost aplikace • u Fermi LI cache (128 byte na řádek) a L2 cache (32 byte na řádek) Lze dynamicky alokovat pomocí cudaMalloc, či staticky pomocí deklarace __c/ew'ce__ JiFÍ Filipovič Úvod, základy CUDA 47/56 Motivace Architektura GPU ooooooooo ooooooo Ostatní paměti Algoritmy a GPU OOOOOOOOO Demonstrační kód oooooooooooo CUDA podrobně ooo«oooooo • paměť konstant • texturová paměť • systémová paměť Uvod, základy CUDA Motivace Architektura GPU ooooooooo ooooooo Algoritmy a GPU OOOOOOOOO Demonstrační kód oooooooooooo CUDA podrobně oooo»ooooo Synchronizace v rámci bloku nativní bariérová synchronizace • musí do ní vstoupit všechny thready (pozor na podmínky!) • pouze jedna instrukce, velmi rychlá, pokud neredukuje paralelismus • v C for CUDA volání __syncthreads() • Fermi rozšíření: count, and, or i -00.0 Uvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooooo oooooo oooooooooooo ooooo«oooo oo Atomické operace • provádí read-modify-write operace nad sdílenou nebo globální pamětí • žádná interference s ostatními thready • pro celá 32-bitová či 64-bitová (pro compute capability > 1.2) čísla (float add u c.c. > 2.0) • nad globální pamětí u zařízení s compute capability > 1.1, nad sdílenou c.c. > 1.2 • aritmetické (Add, Sub, Exch, Min, Max, lne, Dec, CAS) a bitové (And, Or, Xor) operace JiFÍ Filipovič Úvod, základy CUDA 50/56 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooooo oooooo oooooooooooo oooooo«ooo oo Synchronizace paměťových operací Kompilátor může optimalizovat operace se sdílenou/globální pamětí (mezivýsledky mohou zůstat v registrech) a může měnit jejich pořadí, • chceme-li se ujistit, že jsou námi ukládaná data viditelná pro ostatní, používáme __threadfence(), popř. __threadfence_block() • deklarujeme-li proměnnou jako volatile, jsou veškeré přístupy k ní realizovány přes load/store do sdílené či globální paměti • velmi důležité pokud předpokládáme implicitní synchronizaci warpu JiFÍ Filipovič Úvod, základy CUDA 51/56 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooooo oooooo oooooooooooo ooooooo«oo oo Synchronizace bloků Mezi bloky • globální paměť viditelná pro všechny bloky • slabá nativní podpora synchronizace • žádná globální bariéra • u novějších GPU atomické operace nad globální pamětí • globální bariéru lze implementovat voláním kernelu (jiné řešení dosti trikové) • slabé možnosti globální synchronizace znesnadňují programování, ale umožňují velmi dobrou škálovatelnost i -00.0 Uvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooooo oooooo oooooooooooo oooooooo«o oo Globální synchronizace přes atomické operace Problém součtu všech prvků vektoru • každý blok sečte prvky své části vektoru • poslední blok sečte výsledky ze všech bloků • implementuje slabší globální bariéru (po dokončení výpočtu u bloků 1..Í7 — 1 pokračuje pouze blok n) i -00.0 Uvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr OOOOOOOOO OOOOOOO OOOOOOOOO OOOOOO OOOOOOOOOOOO OOOOOOOOO* oo .device__ unsigned int count = 0; .shared__ bool isLastBlockDone; .global__ void sum(const float* array, unsigned int N, float* result) { float partialSum = calculatePartialSum(array, N); if (threadldx.x = 0) { result[blockldx.x] = partialSum; __threadfence(); unsigned int value = at omicine (&count , gridDim.x); isLastBlockDone = (value = (gridDim.x — 1)); } __syncthreads(); if (isLastBlockDone) { float totalSum = calculateTotalSum(result); if (threadldx.x = 0) { result[0] = totalSum; count = 0; } } 4 □ ► 4 fiP ► 4 * -00,0 Uvod. základy CUDA Algoritmy a GPU OOOOOOOOO Demonstrační kód oooooooooooo CUDA podrobně oooooooooo CUDA dokumentace (instalována s CUDA Toolkit, ke stažení na developer.nvidia.com) • CUDA C Programming Guide (nejdůležitější vlastnosti CUDA) • CUDA C Best Practices Guide (detailnější zaměření na optimalizace) • CUDA Reference Manual (kompletní popis C for CUDA API) • další užitečné dokumenty (manuál k nvcc, popis PTX jazyka, manuály knihoven, ...) Série článků CUDA, Supercomputing for the Masses • http: //www.ddj.com/cpp/207200659 4 □ ► 4 fiP ► 4 i -00,0 Uvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooooo oooooo oooooooooooo oooooooooo o« Dnes jsme si ukázali • k čemu je dobré znát CUDA • v čem jsou GPU jiná • základy programování v C for CUDA Příště se zaměříme na • jak psát efektivní GPU kód i -00.0 Uvod, základy CUDA