Architektura GPU Hardwareoua omezeni Demonstrační kod ooooooooooooooooo oooo Obecné výpočty na GPU Jiří Fiiipovič jaro 2010 I—.III Architektura GPU Hardwareoua omezeni Demonstrační kod 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í. I—■M-™" •f) <\(y Architektura GPU Hardwareova omezeni Demonstrační kod 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 ooooooooooooooooo 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 Architektura GPU Hardwareova omezeni Demonstrační kod 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 • 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ů 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í Architektura GPU Hardwareova omezeni Demonstrační kod 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—■M-™" •f) <\(y 11 Architektura GPU Hardwareoua omezeni Demonstrační kod Motivace - výkon GT200 Jan Jun Apr Jun Mař Nov May Jun 2003 2004 2005 2006 2007 2008 □ g - = Architektura GPU Hardwareoua omezeni Demonstrační kod Motivace - výkon Bandwidth GB/s 60 2003 2004 2005 2006 2007 Architektura GPU Hardwareova omezeni Demonstrační kod 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č Architektura GPU Hardwareoua omezeni Demonstrační kod Motivace - uplatnění Využití GPU pro obecné výpočty je dynamicky se rozvíjející oblast s širokou škálou aplikací □ s •f)<\(y \fmmmimiA.M Architektura GPU Hardwareova omezeni Demonstrační kod 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—■M-™" •f) <\(y Architektura GPU Hardwareova omezeni Demonstrační kod 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ódovania dekódování multimediálních dat • herní fyzika • úprava obrázků, 3D rendering • atd... □ gP •f)<\(y 63 Architektura GPU Hardwareova omezeni Demonstrační kod Architektura GPU CPU vs. GPU • jednotky jader vs. desítky m u It i procesorů • out of order vs. in order • MIMD, SIMD pro krátké vektory vs. SIMT pro dlouhé vektory • velká cache vs. malá cache 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í Architektura GPU Architektura GPU CPU dwareova omezeni Demonstrační kod GPU □ ,51 - = ^ ^ q_ o Obecné výpočty na GPU Architektura GPU Hardwareova omezeni Demonstrační kod 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 □ s I—JI! •f) <\(y Architektura GPU Hardwareova omezeni ooooooooooooooooo Demonstrační kod Procesor G80 G80 • první CUDA procesor • obsahuje 16 m u Iti procesorů • m u It i procesor • 8 skalárních procesorů • 2 jednotky pro speciální funkce • až 768 threadů a HW přepínání a plánování threadů • thready organizovány po 32 do warpů • SIMT • nativní synchronizace v rámci multiprocesoru □ s ~ = I—.III •f) <\(y Architektura GPU Hardwareova omezeni Demonstrační kod Paměťový model G80 Paměťový model • 8192 registrů sdílených mezi všemi thready m u Iti procesoru • 16KB sdílené paměti • lokální v rámci m u Iti procesoru » 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 Architektura GPU Hardwareová omezení Demonstrační kód Procesor G80 □ g - = š "O^O Jiří Filipovič Obecné výpočty na GPU ČUDA ČUDA (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 DX11 Compute CUDA Architecture ! Architektura GPU Hardwareova omezeni Demonstrační kod 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 I—JI! •f) <\(y Hierarchie vláken Hierarchie vláken • vlákna jsou organizována do bloků • bloky tvoří mřížku • problém je dekomponovan na podproblemy, které mohou být prováděny nezávisle paralelně (bloky) • jednotlivé podproblemy jsou rozděleny do malých částí, které mohou být prováděny kooperativně paralelně (thready) • dobře ská I uje Motivace Architektura GPU CUDA Hardwareová omezení Demonstrační kód Závěr Hierarchie vláken Grid Block (0,0) Block (1,0) Block (2,0) Block (0,1)- Block (1,1) v Block (2,1) Three Block (1,1) zl Thread (0, 0) Thread (1, 0) Thread (2, 0) Thread (3, 0) Thread (0,1) Thread (1,1) Thread (2,1) Thread (3,1) Thread (1,2) Thread (3, 2) □ &1 - Jiří Filipovič Obecné výpočty na GPU Architektura GPU Hardv Hierarchie pamětí Více druhů pamětí • rozdílná viditelnost • rozdílný čas života • rozdílné rychlosti a chování • přináší dobrou škalovatelnost s ■O Q-C^ ümgMiMJii Architektura GPU Hierarchie pamětí Grid O Block (O, O) Block (1, O) Block (2, O) Block (0,1) Block (1,1) Block (2,1) 1,0) Block (1,0) \fKimmmmAM Architektura GPU Hardwareova omezeni Demonstrační kod Multiprocesor má jen jednu jednotku pro spouštějící instrukce • všech 8 SP musí provádět stejnou instrukci • nová instrukce je spuštěna každé 4 cykly • 32 thredů (tzv. warp) musí provádět stejnou instrukci I—■M-™" •f) <\(y Architektura GPU Hardwareova omezeni Demonstrační kod Multiprocesor má jen jednu jednotku pro spouštějící instrukce • všech 8 SP musí provádět stejnou instrukci • nová instrukce je spuštěna každé 4 cykly • 32 thredů (tzv. warp) musí provádět stejnou instrukci A co větvení kódu? • pokud část threadů ve warpu provádí jinou instrukci, běh se serializuje • to snižuje výkon, snažíme se divergenci v rámci warpu předejít Architektura GPU Hardwareova omezeni Demonstrační kod Multiprocesor má jen jednu jednotku pro spouštějící instrukce • všech 8 SP musí provádět stejnou instrukci • nová instrukce je spuštěna každé 4 cykly • 32 thredů (tzv. warp) musí provádět stejnou instrukci A co větvení kódu? • pokud část threadů ve warpu provádí jinou instrukci, běh se serializuje • to snižuje výkon, snažíme se divergenci v rámci warpu předejít Multiprocesor je tedy MIMD (Multiple-Instruction Multiple-Thread) z programátorského hlediska a SIMT (Single-Instruction Multiple-Thread) z výkonového. Oproti CPU threadům jsou GPU thready velmi lehké (lightweight). • jejich běh může být velmi krátký (desítky instrukcí) • může (mělo by) jich být velmi mnoho • nemohou využívat velké množství prostředků Thready jsou seskupeny v blocích • ty jsou spouštěny na jednotlivých m u Iti procesorech • dostatečný počet bloků je důležitý pro škalovatelnost Počet threadů a thread bloků na multiprocesor je omezen. Architektura GPU Hardwareova omezeni oo»oooooooooooc~ Demonstrační kod Maskování latence pamětí Paměti mají latence • globální paměť má vysokou latenci (stovky cyklů) • registry a sdílená paměť mají read-after-write latenci I—■M-™" •f) <\(y Architektura GPU Hardwareova omezeni oo»oooooooooooc~ Demonstrační kod Maskování latence pamětí Paměti mají latence • globální paměť má vysokou latenci (stovky cyklů) • registry a sdílená paměť mají read-after-write latenci Maskování latence paměti je odlišné, než u CPU • žádné provádění instrukcí mimo pořadí • u většiny pamětí žádná cache □ s Architektura GPU Hardwareova omezeni oo»oooooooooooc~ Demonstrační kod Maskování latence pamětí Paměti mají latence • globální paměť má vysokou latenci (stovky cyklů) • registry a sdílená paměť mají read-after-write latenci Maskování latence paměti je odlišné, než u CPU • žádné provádění instrukcí mimo pořadí • u většiny pamětí žádná cache Když nějaký warp čeká na data z paměti, je možné spustit jiný • umožní maskovat latenci paměti • vyžaduje spuštění řádově více vláken, než má GPU jader • plánování spuštění a přepínání threadů je realizováno přímo v HW bez overhead u Architektura GPU Hardwareova omezeni oo»oooooooooooc~ Demonstrační kod Maskování latence pamětí Paměti mají latence • globální paměť má vysokou latenci (stovky cyklů) • registry a sdílená paměť mají read-after-write latenci Maskování latence paměti je odlišné, než u CPU • žádné provádění instrukcí mimo pořadí • u většiny pamětí žádná cache Když nějaký warp čeká na data z paměti, je možné spustit jiný • umožní maskovat latenci paměti • vyžaduje spuštění řádově více vláken, než má GPU jader • plánování spuštění a přepínání threadů je realizováno přímo v HW bez overhead u Obdobná situace je v případě synchronizace. Architektura GPU Hardwareova omezeni Demonstrační kod Paměťová hierarchie viditelná pro programátora Dekompozice pro GPU • hrubozrnne rozdělení problému na části nevyžadující intenzivní komunikaci/synchronizaci • jemnozrnné rozdělení blízké vektorizaci (SIMT je ale více flexibilní) Bloky mohou využívat sdílenou paměť jako cache. □ s ■O Q-C^ Architektura GPU Hardwareova omezeni oooo»oooooooooooo Demonstrační kod Spojitý přístup do paměti Rychlost GPU paměti je vykoupena nutností přistupovat k ní po větších blocích • globální paměť je dělena do 64-bytových segmentů • ty jsou sdruženy po dvou do 128-bytových segmentů } 64B aligned segment > 128B aligned segment ......I.......... Half warp of threads Architektura GPU Hardwareova omezeni Demonstrační kod Spojitý přístup do paměti Polovina warpu může přenášet data pomocí jedné transakce či jedné až dvou transakcí při přenosu 128-bytového slova • je však zapotřebí využít přenosu velkých slov • jedna paměťová transakce může přenášet 32-, 64-, nebo 128-bytová slova • u GPU sec. < 1.2 • blok paměti, ke kterému je přistupováno, musí začínat na adrese dělitelné šestnáctinásobkem velikosti datových elementů • k-tý thread musí přistupovat ke k-tému elementu bloku • některé thready nemusejí participovat • v případě, že nejsou tato pravidla dodržena, je pro každý element vyvolána zvláštní paměťová transakce GPU s c.c. > 1.2 jsou méně restriktivní • přenos je rozdělen do 32-, 64-, nebo 128-bytových transakcí tak, aby byly uspokojeny všechny požadavky co nejnižším počtem transakcí • pořadí threadů může být vzhledem k přenášeným elementům libovolně permutované Architektura GPU Hardwareoua omezeni ooooooo»ooooooooo Demonstrační kod Spojitý přístup do paměti Thready jsou zarovnané, blok elementů souvislý, pořadí není permutované - spojitý přístup na všech GPU. AA A A A AA A A A—A AAA ET id n ü i—Jíl Architektura GPU Hardwareova omezeni Demonstrační kod 00000»OOOOOOn Nezarovnaný přístup do paměti Thready nejsou zarovnané, blok elementů souvislý, pořadí není permutované -jedna transakce na GPU s c.c. > 1.2. I—.III Architektura GPU Hardwareova omezeni Demonstrační kod ooooooooo«ooooooo oooo Nezarovnaný prístup do paměti Obdobný případ může vézt k nutnosti použít dvě transakce. □ ö - = -š «00,0 Jiří Filipovič Obecné výpočty na GPU Architektura GPU Hardwareoua omezeni oooooooooo«oooooo Demonstrační kod Výkon při nezarovnaném přístupu Starší GPU provádí pro každý element nejmenší možný přenos, tedy 32-bytů, což redukuje výkon na 1/8. Nové GPU (c.c. > 1.2) provádí dva přenosy. -T-GTX28Q f -«-FX5600 D 2 4 B B 10 12 14 16 Offset ■—■M •f) <\(y Architektura GPU Hardwareova omezeni Demonstrační kod ooooooooooo»ooooo oooo Výkon při prokládaném přístupu GPU s c.c. > 1.2 mohou přenášet data s menšími ztrátami pro menší mezery mezi elementy, se zvětšováním mezer výkon dramaticky klesá. D 2 4 B B 10 12 14 16 18 Stride ■—■M = Sdílená paměť je organizována do paměťových bank, ke kterým je možné přistupovat paralelně • 16 bank, paměťový prostor mapován prokládaně s odstupem 32 bitů • pro dosažení plného výkonu paměti musíme přistupovat k datům v rozdílných bankách • implementován broadcast - pokud všichni přistupují ke stejnému údaji v paměti Architektura GPU Hardwareova omezeni ooooooooooooo»ooo Demonstrační kod Konflikty bank Konflikt bank • dojde k němu, přistupují-li některé thready v polovině warpu k datům ve stejné paměťové bance (s výjimkou, kdy všechny thready přistupují ke stejnému místu v paměti) v takovém případě se přístup do paměti serializuje spomalení běhu odpovídá množství paralelních operací, které musí paměť provézt k uspokojení požadavku • je rozdíl, přistupuje-li část threadů k různým datům v jedné bance a ke stejným datům v jedné bance □ ť5> ■O Q-C* Architektura GPU Hardwareova omezeni 0OOOOOOOOOOOOO«O" Demonstrační kod Prístup bez konfliktů □ ► * d5S ► < -= ► OQ,t> Jiří Filipovič Obecné výpočty na GPU Architektura GPU Hardwareova omezeni Demonstrační kod Vřcecestné konflikty 31 ■O Qi O- Jiří Filipovič Obecné výpočty na GPU Architektura GPU Hardwareova omezeni Demonstrační kod ™ÍŤŤľ| VI - s s •e~) q, (y Jiří Filipovič Obecné výpočty na GPU Architektura GPU Hardwareoua omezeni ooooooooooooooooo Demonstrační kod Příklad - součet vektorů Chceme sečíst vektory a a b a výsledek uložit do vektoru c. s ■O Q-C^ ümgMiMJii Architektura GPU Hardwareoua omezeni ooooooooooooooooo Demonstrační kod Příklad - součet vektorů Chceme sečíst vektory aa i?a výsledek uložit do vektoru c. Je třeba najít v problému paralelismus. □ s •f)<\(y \fmmmimiA.M Architektura GPU Hardwareova omezeni ooooooooooooooooo Demonstrační kod 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 < N; i++) □ gP •f)<\(y \fmmmimiA.M Architektura GPU Hardwareova omezeni ooooooooooooooooo Demonstrační kod 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 i < N; b[i]; i++) Jednotlivé iterace cyklu jsou na sobě nezávislé - lze je paralelizovat, škáluje s velikostí vektoru. I—■M-™" •f) <\(y 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? Motivace Architektura GPU C U DA Hardwareová omezení Demonstrační kód Závěr Hierarchie vláken Grid Block (0,0) Block (1,0) Block (2,0) Block (0,1)- Block (1,1) v Block (2,1) Three Block (1,1) zl Thread (0, 0) Thread (1, 0) Thread (2, 0) Thread (3, 0) Thread (0,1) Thread (1,1) Thread (2,1) Thread (3,1) Thread (1,2) Thread (3, 2) □ &1 - Jiří Filipovič Obecné výpočty na GPU C for ČUDA obsahuje zabudované proměnné: • threadldx.jx, y, z} udává pozici threadu v rámci bloku • blockDim.jx, 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) Architektura GPU Hardwareoua omezeni ooooooooooooooooo Demonstrační kod Příklad - součet vektorů Vypočítáme tedy globální pozici threadu (mřížka i bloky jsou jednorozměrné): I—■M-™" •f) <\(y Architektura GPU Hardwareoua omezeni ooooooooooooooooo Demonstrační kod 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 + threadidx.x; I—.III Architektura GPU Hardwareova omezeni ooooooooooooooooo Demonstrační kod Příklad - součet vektorů Vypočítáme tedy globální pozici threadu (mřížka i bloky jsou jednorozměrné): blockIdx.x*blockDin 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]; I—JI! •f) <\(y 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 threadu a v jakém uspořádání bude spuštěno. Jaké výpočty urychlovat? • velké instance problémů • kritické pro výkon aplikace • možná dekompozice do tisíců threadů • žádná nebo málo častá globální synchronizace • vektorový charakter výpočtů • vysoký podíl aritmetických operací, nebo dostatek paměťových operací na vstupních datech □ s Zhodnocení GPU představují výkonné, běžně dostupné akcelerátory • řádový posun výkonu je dostatečně zajímavý • lze předpokládat, že se jejich náskok bude dále zvyšovat • jsou levné a rozšířené Programovací model je složitý, ale zvládnutelný • z programování CPU nejsme zvyklí na tak restriktivní výkonová omezení » použitelnost programovacího modelu dokazují úspěšné aplikace z mnoha oblastí <» stejně se paralelní programování musíme naučit :-) • otevřené pole jak pro vývoj aplikací, tak pro výzkum □ s Architektura GPU Hard Pokud vás téma zaujalo. Kam dál na Fl? • PV197 GPU Programming • napište na fila@mail.muni.cz s ■O Q-C^ ümgMiMJii