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 2012 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 Moorův zákon Počet tranzistorů na jednom čipu se přibližně každých 18 měsíců zdvojnásobí. Jih Filipovič Úvod, 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 Moorů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 Jih Filipovič Úvod, 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 Jih Filipovič Úvod, základy CUDA 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ů Jih Filipovič Úvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr 000*00000 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í Jih Filipovič Úvod, 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 Jih Filipovič Úvod, 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 Jih Filipovič Úvod, základy CUDA 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 Jiří Filipovič Úvod, základy CUDA 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č Jih Filipovič Úvod, 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č Úvod, 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ší... Jih Filipovič Úvod. 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... Jih Filipovič Úvod. základy CUDA 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í Jih Filipovič Úvod. 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 Jih Filipovič Úvod, základy CUDA 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 Jih Filipovič Úvod, 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 Jih Filipovič Úvod, základy CUDA 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 Jiří Filipovič Úvod, základy CUDA 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 Jih Filipovič Úvod, základy CUDA 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> Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr OOOOOOOOO OOOOOOO »00000000 oooooo oooooooooooo oooooooooo oo 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 poukazuje na hloupost autorů Srovnáváme-li přínos GPU oproti CPU, musíme uvažovat efektivní implementaci pro obě platformy. Jiří Filipovič Úvod, základy CUDA 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í. Jiří Filipovič Úvod, základy CUDA 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. Jiří Filipovič Úvod, základy CUDA 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 Jih Filipovič Úvod, základy CUDA 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) Jih Filipovič Úvod, základy CUDA 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 (!) Jiří Filipovič Úvod, základy CUDA 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 Jih Filipovič Úvod, základy CUDA 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 Jih Filipovič Úvod, základy CUDA 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 Jih Filipovič Úvod, 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++ DX11 Compute CUDA Architecture i -00,0 Jiří Filipovič Úvod, 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 Jih Filipovič Úvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooooo oocooo 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 Jih Filipovič Úvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr OOOOOOOOO OOOOOOO OOOOOOOOO OOOÍOO oooooooooooo oooooooooo oo Hierarchie vláken Grid Block (0,0) 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) Jiří Filipovič Úvod, základy CUDA 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 Jih Filipovič Úvod, 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í Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU ooooooooo ooooooo ooooooooo CUDA Demonstrační kód OOOOOO «00000000000 CUDA podrobně oooooooooo Závěr 00 Příklad - součet vektorů Chceme sečíst vektory a a b a výsledek uložit do vektoru c. Jih Filipovič Úvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU ooooooooo ooooooo ooooooooo CUDA Demonstrační kód OOOOOO «00000000000 CUDA podrobně oooooooooo Závěr 00 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. Jih Filipovič Úvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU ooooooooo ooooooo ooooooooo CUDA Demonstrační kód OOOOOO «00000000000 CUDA podrobně oooooooooo Závěr 00 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]; Jih Filipovič Úvod. základy CUDA Motivace Architektura GPU Algoritmy a GPU ooooooooo ooooooo ooooooooo CUDA Demonstrační kód OOOOOO «00000000000 CUDA podrobně oooooooooo Závěr 00 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. Jih Filipovič Úvod. základy CUDA Motivace Architektura GPU Algoritmy a GPU ooooooooo ooooooo ooooooooo CUDA Demonstrační kód OOOOOO «00000000000 CUDA podrobně oooooooooo Závěr 00 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? Jih Filipovič Úvod. základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooooo oooooo o»oooooooooo oooooooooo oo Hierarchie vláken Grid Block (0,0) 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) Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooooo oooooo oo«ooooooooo oooooooooo oo 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.(x, 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) Jih Filipovič Úvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooooo oooooo ooo»oooooooo oooooooooo oo Příklad - součet vektorů Vypočítáme tedy globální pozici threadu (mřížka i bloky jsou jednorozměrné): Jih Filipovič Úvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooooo oooooo ooo»oooooooo oooooooooo oo 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; Jih Filipovič Úvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooooo oooooo ooo»oooooooo oooooooooo oo 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; 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]; } Jih Filipovič Úvod. základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooooo oooooo ooo»oooooooo oooooooooo oo 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; 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. Jih Filipovič Úvod. základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooooo oooooo oooo«ooooooo oooooooooo oo 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í Jih Filipovič Ú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: Jih Filipovič Ú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 Jih Filipovič Ú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 Jih Filipovič Ú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 Jih Filipovič Ú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 Jih Filipovič Ú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 • uložit výsledek z GPU paměti do c Jih Filipovič Ú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 • uložit výsledek z GPU paměti do c • použít výsledek v c :-) Jih Filipovič Úvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU ooooooooo ooooooo ooooooooo CUDA Demonstrační kód OOOOOO 000000*00000 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++) a[i] = b[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; } Jih Filipovič Úvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU ooooooooo ooooooo ooooooooo CUDA Demonstrační kód oooooo ooooooo«oooo CUDA podrobně oooooooooo Závěr 00 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). Jih Filipovič Úvod. 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); Jih Filipovič Úvod, základy CUDA 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? Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooooo oooooo oooooooooo»o oooooooooo oo 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); Jih Filipovič Úvod, 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 hi-end GPU: barracuda.fi.muni.cz, airacuda.fi.muni.cz, účty na přání Jih Filipovič Úvod. základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA ooooooooo ooooooo ooooooooo oooooo Demonstrační kód oooooooooooo CUDA podrobně •OOOOOOOOO Závěr 00 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) Jih Filipovič Úvod, 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 nevleze 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) Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooooo oooooo oooooooooooo o«oooooooo oo 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 Jih Filipovič Úvod. 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__ Jih Filipovič Úvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr OOOOOOOOO OOOOOOO OOOOOOOOO oooooo oooooooooooo ooo«oooooo oo Ostatní paměti • paměť konstant • texturová paměť • systémová paměť Jih Filipovič Úvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Závěr ooooooooo ooooooo ooooooooo oooooo oooooooooooo oooo»ooooo oo 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 Jih Filipovič Úvod, 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 Jih Filipovič Úvod. základy CUDA 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 Jih Filipovič Úvod, základy CUDA 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 Jih Filipovič Úvod, základy CUDA Motivace Architektura GPU Algoritmy a GPU ooooooooo ooooooo ooooooooo Globální synchronizace přes CUDA Demonstrační kód CUDA podrobně Závěr OOOOOO OOOOOOOOOOOO 00000000*0 OO 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) Jih Filipovič Úvod, 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; } } } Jiří Filipovič Úvod. 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 Materiály 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, ...) Textbook ke kurzům na University of Illinois • dostupný z http://courses.ece. i llinois.edu/ece498/al/Syl la bus. html Série článků CUDA, Supercomputing for the Masses • http: //www.ddj.com/cpp/207200659 Jiří Filipovič Úvod, 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 Jih Filipovič Úvod, základy CUDA