Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo ooooooooo oo oooooooooooooooooo Optimalizace pro GPU hardware Jiří Filipovič jaro 2015 Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce •ooo oooooooooooooooooo ooooooooo oo oooooooooooooooooo Paralelismus GPU Paralelní algoritmy je nutno navrhovat vzhledem k paralelismu, který poskytuje HW • v případě GPU se jedná o pole SIMT m u Iti procesorů pracující nad společnou pamětí Dekompozice pro GPU • hrubozrnné rozdělení problému na části nevyžadující intenzivní komunikaci/synchronizaci • jemnozrnné rozdělení blízké vektorizaci (SIMT je ale více flexibilní) Jiří Filipovič Optimalizace pro GPU hardware 2/52 Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOO Transpozice matic OOOOOOOOO Rychlost instrukcí 00 Na jednu jednotku spouštějící instrukce připadá několik skalárních procesorů (SP) G80 • 8 SP na jednotku spouštějící instrukce • nová instrukce je spuštěna každé 4 cykly • 32 thredů (tzv. warp) musí provádět stejnou instrukci * si ► * Optimalizace pro GPU hardware i -00.0 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce o»oo OOOOOOOOOOOOOOOOOO OOOOOOOOO OO OOOOOOOOOOOOOOOOOO SIMT Na jednu jednotku spouštějící instrukce připadá několik skalárních procesorů (SP) G80 • 8 SP na jednotku spouštějící instrukce • 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 Jiří Filipovič Optimalizace pro GPU hardware 3/52 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce 0«00 OOOOOOOOOOOOOOOOOO OOOOOOOOO OO OOOOOOOOOOOOOOOOOO SIMT Na jednu jednotku spouštějící instrukce připadá několik skalárních procesorů (SP) G80 • 8 SP na jednotku spouštějící instrukce • 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. Jiří Filipovič Optimalizace pro GPU hardware 3/52 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oo«o oooooooooooooooooo ooooooooo oo oooooooooooooooooo Vlastnosti threadů Oproti CPU threadům jsou GPU thready velmi „jemné". • 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 multiprocesorech • dostatečný počet bloků je důležitý pro škálovatelnost Počet threadů a thread bloků na multiprocesor je omezen. Jiří Filipovič Optimalizace pro GPU hardware 4/52 Optimalizace přístupu do paměti oooooooooooooooooo Transpozice matic ooooooooo Rychlost instrukcí 00 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 Optimalizace pro GPU hardware i -00.0 Optimalizace přístupu do paměti oooooooooooooooooo Transpozice matic ooooooooo Rychlost instrukcí 00 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í • často žádná cache Optimalizace pro GPU hardware i -00.0 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce ooo» oooooooooooooooooo ooooooooo oo oooooooooooooooooo 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í • často žá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 overheadu Jiří Filipovič Optimalizace pro GPU hardware 5/52 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce ooo» oooooooooooooooooo ooooooooo oo oooooooooooooooooo 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í • často žá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 overheadu Obdobná situace je v případě synchronizace. Jiří Filipovič Optimalizace pro GPU hardware 5/52 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO «00000000000000000 OOOOOOOOO OO OOOOOOOOOOOOOOOOOO Optimalizace přístupu do globální paměti Rychlost globální paměti se snadno stane bottleneckem • šířka pásma globální paměti je ve srovnání s aritmetickým výkonem GPU malá (desítky flops na přenos slova) • latence 400-600 cyklů Při špatném vzoru paralelního přístupu do globální paměti snadno výrazně snížíme propustnost • k paměti je nutno přistupovat sdruženě (coalescing) 9 je vhodné vyhnout se užívání pouze podmnožiny paměťových regionů (partition camping) Jiří Filipovič Optimalizace pro GPU hardware 6/52 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO OOOOOOOOOOOOOOOOOO OOOOOOOOO OO OOOOOOOOOOOOOOOOOO Sdružený přístup do paměti (c.c. < 2.0) 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 J128B aligned segment ................. Half warp of threads Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO oo«ooooooooooooooo OOOOOOOOO OO OOOOOOOOOOOOOOOOOO Sdružený přístup do paměti (c.c. < 2.0) 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 s c.c. < 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 Jiří Filipovič Optimalizace pro GPU hardware 8/52 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO OOO^OOOOOOOOOOOOOO OOOOOOOOO OO OOOOOOOOOOOOOOOOOO Sdružený přístup do paměti (c.c. < 2.0) 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 nej nižším počtem transakcí • pořadí threadů může být vzhledem k přenášeným elementům libovolně permutované Optimalizace pro GPU hardware i -00.0 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO OOOOOOOOOOOOOOOOOO OOOOOOOOO OO OOOOOOOOOOOOOOOOOO Sdružený přístup do paměti (c.c. < 2.0) Thready jsou zarovnané, blok elementů souvislý, pořadí není permutované - sdružený přístup na všech GPU. á i J 1 1 I 1 1 1 1 r t 1 1 1 1 Optimalizace pro GPU hardware i -00.0 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo ooooooooo oo oooooooooooooooooo Nezarovnaný přístup do paměti (c.c. < 2.0) Thready nejsou zarovnané, blok elementů souvislý, pořadí není permutované - jedna transakce na GPU s c.c. > 1.2. 'ľľľľľľl Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO OOOOOOOOOOOOOOOOOO OOOOOOOOO OO OOOOOOOOOOOOOOOOOO Nezarovnaný přístup do paměti(c.c. < 2.0) Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo OOOOOOOOOOOOOOOOOO ooooooooo oo oooooooooooooooooo Výkon při nezarovnaném přístupu (c.c. < 2.0) 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. -W-GTX280 ? -*FX5600 2 4 6 8 10 12 14 16 Offset Optimalizace pro GPU hardware i -00.0 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO 00000000«000000000 OOOOOOOOO OO OOOOOOOOOOOOOOOOOO Výkon při prokládaném přístupu (c.c. < 2.0) 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á. 0 2 4 6 8 10 12 14 16 18 Stride Jiří Filipovič Optimalizace pro GPU hardware 14 /52 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO 000000000*00000000 OOOOOOOOO OO OOOOOOOOOOOOOOOOOO Přístup do globální paměti u Fermi (c.c. > 2.0) Fermi má LI a L2 cache • LI: 256 byte na řádek, celkem 16 KB nebo 48 KB na m u Iti procesor • L2: 32 byte na řádek, celkem 768 KB na GPU Jaké to přináší výhody? • efektivnější programy s nepředvídatelnou datovou lokalitou • nezarovnaný přístup - v principu žádné spomalení » prokládaný přístup - data musí být využita dříve, než zmizí z cache, jinak stejný či větší problém jako u c.c. < 2.0 (LI lze vypnout pro zamezení overfetchingu) Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO 0000000000*0000000 ooooooooo oo oooooooooooooooooo HW organizace sdílené paměti Sdílená paměť je organizována do paměťových bank, ke kterým je možné přistupovat paralelně • c.c. 1.x 16 bank, c.c. 2.x 32 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 Jiří Filipovič Optimalizace pro GPU hardware 16 /52 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo ooooooooooo«oooooo ooooooooo oo oooooooooooooooooo Konflikty bank Konflikt bank • dojde k němu, přistupují-li některé thready v warpu/půlwarpu k datům ve stejné paměťové bance (s výjimkou, kdy 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ést 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 Jiří Filipovič Optimalizace pro GPU hardware 17 /52 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooo»ooooo ooooooooo oo oooooooooooooooooo Přístup bez konfliktů Jiří Filipovič Optimalizace pro GPU hardware 18 /52 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO OOOOOOOOOOOOOOOOOO OOOOOOOOO OO OOOOOOOOOOOOOOOOOO Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO OOOOOOOOOOOOOO^OOO OOOOOOOOO OO OOOOOOOOOOOOOOOOOO Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO ooooooooooooooo«oo OOOOOOOOO OO OOOOOOOOOOOOOOOOOO Vzory přístupu Zarovnání není třeba, negeneruje bank conflicty int x = s[threadldx.x + offset]; Prokládání negeneruje konflikty, je-li c liché int x = s[threadldx.x * c]; Přístup ke stejné proměnné negeneruje na c.c. 2.x konflikty nikdy, na l,x je-li počet c threadů přistupující k proměnné násobek 16 int x = s[threadldx.x / c]; Optimalizace pro GPU hardware i -00.0 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO OOOOOOOOOOOOOOOOOO OOOOOOOOO OO OOOOOOOOOOOOOOOOOO Ostatní paměti Přenosy mezi systémovou a grafickou pamětí • je nutné je minimalizovat (často i za cenu neefektivní části výpočtu na GPU) • mohou být zrychleny pomcí page-locked paměti • je výhodné přenášet větší kusy současně • je výhodné překrýt výpočet s přenosem Texturová paměť • vhodná k redukci přenosů z globální paměti • vhodná k zajištění sdruženého přístupu • nevhodná, pokud je bottleneck latence • může zjednodušit adresování či přidat filtraci Jiří Filipovič Optimalizace pro GPU hardware 22 /52 Paralelismus OOOO Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOO* Transpozice matic OOOOOOOOO Rychlost instrukcí 00 Ostatní paměti Paměť konstant • rychlá jako registry, pokud čteme tutéž hodnotu • se čtením různých hodnot lineárně klesá rychlost Registry • read-after-write latence, odstíněna pokud na m u Iti procesoru běží alespoň 192 threadů pro c.c. 1.x a 768 threadů pro c.c. 2.x • potenciální bank konflikty i v registrech • kompilátor se jim snaží zabránit • můžeme mu to usnadnit, pokud nastavíme velikost bloku na násobek 64 Optimalizace pro GPU hardware i -00.0 Paralelismus OOOO Optimalizace přístupu do paměti oooooooooooooooooo Transpozice matic ♦OOOOOOOO Rychlost instrukcí 00 Transpozice matic Z teoretického hlediska: • triviální problém • triviální paralelizace • jsme triviálně omezení propustností paměti (neděláme žádné flops) __global__ void mtran(float *odata, float* idata, int n){ int x = blockldx.x * blockDim.x + threadldx.x; int y = blockldx.y * blockDim.y + threadldx.y; odata[x*n + y] = idata[y*n + x]; Optimalizace pro GPU hardware i -00.0 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO OOOOOOOOOOOOOOOOOO OÄOOOOOOO OO OOOOOOOOOOOOOOOOOO Výkon Spustíme-li kód na GeForce GTX 280 s použitím dostatečně velké matice 4000 x 4000, bude propustnost 5.3 GB/s. Kde je problém? Jiří Filipovič Optimalizace pro GPU hardware 25 /52 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO OOOOOOOOOOOOOOOOOO OÄOOOOOOO OO OOOOOOOOOOOOOOOOOO Výkon Spustíme-li kód na GeForce GTX 280 s použitím dostatečně velké matice 4000 x 4000, bude propustnost 5.3 GB/s. Kde je problém? Přístup do odata je prokládaný! Modifikujeme transpozici na kopírování: odata[y*n + x] = idata[y*n + x]; a získáme propustnost 112.4 GB/s. Pokud bychom přistupovali s prokládáním i k idata, bude výsledná rychlost 2.7 GB/s. Jiří Filipovič Optimalizace pro GPU hardware 25 /52 Paralelismus OOOO Optimalizace přístupu do paměti oooooooooooooooooo Transpozice matic oo«oooooo Rychlost instrukcí 00 Odstranění prokládání Matici můžeme zpracovávat po dlaždicích • načteme po řádcích dlaždici do sdílené paměti • uložíme do globální paměti její transpozici taktéž po řádcích • díky tomu je jak čtení, tak zápis bez prokládání * si ► * Optimalizace pro GPU hardware i -00.0 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo oo«oooooo oo oooooooooooooooooo Odstranění prokládání Matici můžeme zpracovávat po dlaždicích • načteme po řádcích dlaždici do sdílené paměti • uložíme do globální paměti její transpozici taktéž po řádcích • díky tomu je jak čtení, tak zápis bez prokládání Jak velké dlaždice použít? » budeme uvažovat dlaždice čtvercové velikosti • pro sdružené čtení musí mít řádek dlaždice velikost dělitelnou 16 • v úvahu připadají dlaždice 16 x 16, 32 x 32 a 48 x 48 (jsme omezeni velikostí sdílené paměti) • nejvhodnější velikost určíme experimentálně Jiří Filipovič Optimalizace pro GPU hardware 26 /52 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo ooo»ooooo oo oooooooooooooooooo Dlaždicová transpozice __global__ void mtran_coalesced(float *odata, float *idata, int n) __shared__ float tile[TILE_DIM][TILE_DIM]; int x — blockldx.x * TILE_DIM + threadldx.x; int y — blockldx.y * TILE_DIM + threadldx.y; int index_in = x + y*n; x — blockldx.y * TILE_DIM + threadldx.x; y — blockldx.x * TILE_DIM + threadldx.y; int index_out = x + y*n ; for (int i = 0; i < TILE_DIM; i += BL0CK_R0WS) tile[threadldx.y+i][threadldx.x] = idata[index.in+i*n]; __syncthreads(); for (int i = 0; i < TILE_DIM; i += BL0CK_R0WS) odata[index_out+i*n] = tile[threadldx.x][threadldx.y+i]; } Jiří Filipovič Optimalizace pro GPU hardware 27 /52 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO OOOOOOOOOOOOOOOOOO oooo«oooo OO OOOOOOOOOOOOOOOOOO Výkon Nejvyšší výkon byl naměřen při použití dlaždic velikosti 32 x 32, velikost thread bloku 32 x 8, a to 75.1GB/s. Jiří Filipovič Optimalizace pro GPU hardware 28 /52 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo oooo«oooo oo oooooooooooooooooo Výkon Nejvyšší výkon byl naměřen při použití dlaždic velikosti 32 x 32, velikost thread bloku 32 x 8, a to 75.1GB/s. • to je výrazně lepší výsledek, nicméně stále nedosahujeme rychlosti pouhého kopírování Optimalizace pro GPU hardware i -00.0 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo oooo«oooo oo oooooooooooooooooo Výkon Nejvyšší výkon byl naměřen při použití dlaždic velikosti 32 x 32, velikost thread bloku 32 x 8, a to 75.1GB/s. • to je výrazně lepší výsledek, nicméně stále nedosahujeme rychlosti pouhého kopírování • kernel je však složitější, obsahuje synchronizaci • je nutno ověřit, jestli jsme narazili na maximum, nebo je ještě někde problém Optimalizace pro GPU hardware i -00.0 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO oooooooooooooooooo oooo«oooo OO OOOOOOOOOOOOOOOOOO Výkon Nejvyšší výkon byl naměřen při použití dlaždic velikosti 32 x 32, velikost thread bloku 32 x 8, a to 75.1GB/s. • to je výrazně lepší výsledek, nicméně stále nedosahujeme rychlosti pouhého kopírování • kernel je však složitější, obsahuje synchronizaci • je nutno ověřit, jestli jsme narazili na maximum, nebo je ještě někde problém • pokud pouze kopírujeme, dosáhneme výkonu 94.9GB/s • něco ještě není optimální Jiří Filipovič Optimalizace pro GPU hardware 28 /52 Paralelismus OOOO Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOO Transpozice matic ooooo»ooo Rychlost instrukcí 00 Sdílená paměť Při čtení globální paměti zapisujeme do sdílené paměti po řádcích. t ile[threadldx.y+i] [threadldx. x ] = idata[index_ in+i*n]; Optimalizace pro GPU hardware i -00.0 Paralelismus OOOO Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOO Transpozice matic OOOOO0OOO Rychlost instrukcí 00 Sdílená paměť Při čtení globální paměti zapisujeme do sdílené paměti po řádcích. t ile[threadldx.y+i] [threadldx.x] = idata[index_in+i*n]; Při zápisu do globální paměti čteme ze sdílené po sloupcích. odata[index_out+i*n] = tile[threadldx.x][threadldx.y+i]; To je čtení s prokládáním, které je násobkem 16, celý sloupec je tedy v jedné bance, vzniká 16-cestný bank conflict. Optimalizace pro GPU hardware i -00.0 Paralelismus OOOO Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOO Transpozice matic ooooo»ooo Rychlost instrukcí 00 Sdílená paměť Při čtení globální paměti zapisujeme do sdílené paměti po řádcích. tile[threadldx.y+i][threadldx.x] = idata[index.in+i*n]; Při zápisu do globální paměti čteme ze sdílené po sloupcích. odata[index_out+i*n] = tile[threadldx.x][threadldx.y+i] ; To je čtení s prokládáním, které je násobkem 16, celý sloupec je tedy v jedné bance, vzniká 16-cestný bank conflict. Řešením je padding: __shared__ float tile[TILE_DIM][TILE_DIM + 1]; Optimalizace pro GPU hardware i -00.0 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO OOOOOOOOOOOOOOOOOO OOOOOOÄOO OO OOOOOOOOOOOOOOOOOO Výkon Nyní dosahuje naše implementace výkon 93.4 GB/s. • obdobný výsledek, jako při pouhém kopírování • zdá se, že výrazněji lepšího výsledku již pro danou matici nedosáhneme • pozor na různou velikost vstupních dat (tzv. partition camping, není problém u Fermi) Jiří Filipovič Optimalizace pro GPU hardware 30 / 52 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo ooooooo»o oo oooooooooooooooooo Zhodnocení výkonu Veškeré optimalizace sloužily pouze k lepšímu přizpůsobení-se vlastnostem HW • přesto jsme dosáhli 17.6x zrychlení • při formulaci algoritmu je nezbytné věnovat pozornost hardwareovým omezením • jinak bychom se nemuseli vývojem pro GPU vůbec zabývat, stačilo by napsat dobrý CPU algoritmus... Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO OOOOOOOOOOOOOOOOOO 00000000« oo oooooooooooooooooo Význam optimalizací Pozor na význam optimalizací • pokud bychom si zvolili testovací matice velikosti 4096 x 4096 namísto 4000 x 4000, byl by efekt odstranění konfliktů ve sdílené paměti po odstranění prokládaného přístupu prakticky neznatelný • po odstranění partition campingu by se však již konflikty bank výkonnostně projevily! • je dobré postupovat od obecně zásadnějších optimalizací k těm méně zásadním • nevede-li některá (prokazatelně korektní :-)) optimalizace ke zvýšení výkonu, je třeba prověřit, co algoritmus brzdí Jiří Filipovič Optimalizace pro GPU hardware 32 /52 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo ooooooooo «o oooooooooooooooooo Rychlost provádění instrukcí Některé instrukce jsou v porovnání s ostatními pomalejší, než u procesoru • celočíselné dělení a modulo • 32-bitové násobení u c.c. 1.x • 24-bitové násobení u c.c. 2.x Některé jsou naopak rychlejší • méně přesné verze prováděné na SFU • sinf(x), __cosf(x), __expf(x), __sincosf(x), __rsqrtf(x) aj. Jiří Filipovič Optimalizace pro GPU hardware 33 / 52 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO OOOOOOOOOOOOOOOOOO OOOOOOOOO 0» OOOOOOOOOOOOOOOOOO Smyčky Malé cykly mají značný overhead • je třeba provádět skoky • je třeba updatovat kontrolní proměnnou • podstatnou část instrukcí může tvořit pointerová aritmetika To lze řešit rozvinutím (unrolling) • částečně je schopen dělat kompilátor • můžeme mu pomoci ručním unrollingem, nebo pomocí direktivy #pragma unroll Jiří Filipovič Optimalizace pro GPU hardware 34 /52 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO OOOOOOOOOOOOOOOOOO OOOOOOOOO OO »00000000000000000 Součet prvků vektoru Pro vektor v o n prvcích chceme spočítat x = Y^i=i v'f- Jiří Filipovič Optimalizace pro GPU hardware 35 /52 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO OOOOOOOOOOOOOOOOOO OOOOOOOOO OO »00000000000000000 Součet prvků vektoru Pro vektor v o n prvcích chceme spočítat x = Y11=i vi-Zápis v jazyce C int x = O; for (int i — 0; i < n; i++) x += v[i ] ; Jednotlivé iterace cyklu jsou na sobě závislé. Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO OOOOOOOOOOOOOOOOOO OOOOOOOOO OO »00000000000000000 Součet prvků vektoru Pro vektor v o n prvcích chceme spočítat x = Y11=i vi-Zápis v jazyce C int x = O; for (int i — 0; i < n; i++) x += v[i ] ; Jednotlivé iterace cyklu jsou na sobě závislé. • nemůžeme udělat všechnu práci paralelně • sčítání je však (alespoň teoreticky :-)) asocitativní • není tedy nutno počítat sekvenčně Optimalizace pro GPU hardware Optimalizace přístupu do paměti oooooooooooooooooo Transpozice matic ooooooooo Rychlost instrukcí 00 Paralelní algoritmus Představený sekvenční algoritmus provádí pro 8 prvků výpočet: ((((((1/1 + v2) + v3) + 1/4) + 1/5) + v6) + vr) + V8 Optimalizace pro GPU hardware Optimalizace přístupu do paměti oooooooooooooooooo Transpozice matic ooooooooo Rychlost instrukcí 00 Paralelní algoritmus Představený sekvenční algoritmus provádí pro 8 prvků výpočet: ((((((1/1 + v2) + 1/3) + 1/4) + 1/5) + v6) + vr) + V8 Sčítání je asociativní... spřeházejme tedy závorky: {{vi + v2) + (1/3 + 1/4)) + ((1/5 + v6) + {vj + v8)) Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo ooooooooo oo o«oooooooooooooooo Paralelní algoritmus Představený sekvenční algoritmus provádí pro 8 prvků výpočet: ((((((1/1 + v2) + v3) + 1/4) + 1/5) + v6) + vr) + v8 Sčítání je asociativní... spřeházejme tedy závorky: {{vi + v2) + (1/3 + 1/4)) + ((1/5 + v6) + {vj + v8)) Nyní můžeme pracovat paralelně • v prvním kroku provedeme 4 sčítání • ve druhém dvě • ve třetím jedno Celkově stejné množství práce (n — 1 sčítání), ale v log2 n paralelních krocích! Jiří Filipovič Optimalizace pro GPU hardware 36 /52 Optimalizace přístupu do paměti oooooooooooooooooo Transpozice matic ooooooooo Rychlost instrukcí 00 Paralelní algoritmus Našli jsme vhodný paralelní algoritmus • provádí stejné množství operací jako sériová verze • při dostatku procesorů je proveden v logaritmickém čase Sčítáme výsledky předešlých součtů • předešlé součty provádělo více threadů • vyžaduje globální bariéru Optimalizace pro GPU hardware i -00.0 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO OOOOOOOOOOOOOOOOOO OOOOOOOOO OO ooo«oooooooooooooo Naivní přístup Nejjednodušší schéma algoritmu: • kernel pro sudá i < n provede v[i] += v[i+l] • opakujeme pro n /= 2 dokud n > 1 Omezení výkonu • 2n čtení z globální paměti • n zápisů do globální paměti • log2 n volání kernelu Na jednu aritmetickou operaci připadají 3 paměťové přenosy, navíc je nepříjemný overhead spouštění kernelu. Jiří Filipovič Optimalizace pro GPU hardware 38 / 52 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo ooooooooo oo oooo»ooooooooooooo Využití rychlejší paměti V rámci volání kernelu můžeme posčítat více, než jen dvojice • každý blok bx načte m prvků do sdílené paměti • provede redukci (ve sdílené paměti v log2 m krocích) • uloží pouze jedno číslo odpovídající Yl!h=bmXx y> Výhodnější z hlediska paměťových přenosů i spouštění kernelů •" + ^ + ^ + •• + ^ = (""1)^1 • přibližně n + ^ čtení, ^ zápisů • logm n spuštění kernelu Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo ooooooooo oo ooooo«oooooooooooo Implementace 1 .global__ void reducel(int *v){ extern__shared__ int sv [1 ; unsigned int tid = threadldx.x; unsigned int i = blockldx.x*blockDim.x + threadldx.x; sv [ t id ] = v [ i ] ; __syncthreads (); for(unsigned int s=l; s < blockDim.x; s *= if (tid % (2*s) == 0) s v [tid] += s v [ t i d + s]; __syncthreads(); } 2) { if (tid = 0) v[blockldx.x] = sv[0]; i -00.0 Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO OOOOOOOOOOOOOOOOOO OOOOOOOOO OO OOOOOOOOOOOOOOOOOO Výkon Vysoká úroveň divergence • první iteraci pracuje každý 2. thread • druhou iteraci pracuje každý 4. thread • třetí iteraci pracuje každý 8 thread • atd. Přenos (GTX 280) 3.77GB/s, 0.94MEIem/s. Jiří Filipovič Optimalizace pro GPU hardware 41 /52 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo ooooooooo oo ooooooo«oooooooooo Implementace 2 Nahradíme indexaci ve for cyklu for (unsigned int s = 1; s < blockDim.x; int index = 2 * s * tid; if (index < blockDim.x) sv[index] += sv[index + s]; __syncthreads(); } Přenos 8.33GB/s, 2.08MEIem/s. Řeší divergenci, generuje konflikty bank. 2) { Optimalizace pro GPU hardware i -00.0 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo ooooooooo oo oooooooo«ooooooooo Implementace 3 Tak ještě jinak... for (unsigned int s — blockDim . x/2; s > 0; s »= 1) { if (tid < s) s v [tid] += s v[t i d + s]; __syncthreads () ; } Žádná divergence ani konflikty. Přenos 16.34 GB/s, 4.08MEIem/s. Polovina threadů nic nepočítá... Optimalizace pro GPU hardware i -00.0 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo ooooooooo oo ooooooooo«oooooooo Implementace 4 První sčítání provedeme již během načítání. unsigned int i sv[t id 1 = v [ i I blockldx.x*(blockDim.x * 2) v[i+blockDim. x 1 ; threadldx.x; Přenos 27.16 GB/s, 6.79MEIem/s. Data zřejmě čteme optimálně, stále je zde však výkonová rezerva zaměřme se na instrukce. * si ► * Optimalizace pro GPU hardware i -00.0 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo ooooooooo oo oooooooooo«ooooooo Implementace 5 V jednotlivých krocích redukce ubývá aktivních threadů • nakonec bude pracovat pouze jeden warp • ten je však synchronizován implicitně, můžeme tedy odebrat syncthreadsQ • podmínka if[tid < s) je zde zbytečná (nic neušetří) Unrollujme tedy poslední warp... * si ► * Optimalizace pro GPU hardware i -00.0 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO OOOOOOOOOOOOOOOOOO OOOOOOOOO OO OOOOOOOOOOOÍOOOOOO Implementace 5 for (unsigned int s = blockDim . x/2; s > 32; s »= 1){ if (tid < s) s v [tid] += s v[t i d + s]; __syncthreads () ; } if (tid < 32){ s v tid] += sv tid 4 - 32] s v tid] += sv tid 4 - 16] s v tid] += sv tid 4 - 8]; s v tid] += sv tid 4 - 4]; s v tid] sv tid 4 - 2]; s v tid] += sv tid 4 - i]; } Ušetříme čas i ostatním waprům (skončí dříve s for cyklem). Přenos 37.68 GB/s, 9.42MEIem/s. Jiří Filipovič Optimalizace pro GPU hardware 46 /52 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo ooooooooo oo oooooooooooo«ooooo Implementace 6 Jak je to s rozvinutím for cyklu? Známe-li počet iterací, můžeme cyklus rozvinout • počet iterací je závislý na velikosti bloku Můžeme být obecní? • algoritmus pracuje s bloky o velikosti 2" • velikost bloku je shora omezena • známe-li při kompilaci velikost bloku, můžeme použít šablonu template __global__ void reduce6(int *v) Jiří Filipovič Optimalizace pro GPU hardware 47 /52 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO OOOOOOOOOOOOOOOOOO OOOOOOOOO OO ooooooooooooo«oooo Podmínky s blockSize se vyhodnotí již pří překladu: if (blockSize >= 512){ if (tid < 256) sv[tid] += sv[tid + 256]; __syncthreads () ; } if (blockSize >= 256){ if (tid < 128) s v[tid] += s v[t i d + __syncthreads () ; 128]; } if (blockSize >= 128){ if (tid < 64) s v[tid] += s v[t i d + __syncthreads ( ) ; 64]; } Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo ooooooooo oo oooooooooooooo«ooo Implementace 6 if (tid < 32){ if (blocksize >= 64) s v[tid] += s v[tid 4- 32] if (blocksize >= 32) s v[tid] += s v[tid 4- 16] if (blocksize >= 16) s v[tid] += s v[tid 4- 8]; if (blocksize >= 8) s v[tid] 4-= sv [ tid + 4]; if (blocksize >= 4) s v[tid] 4-= sv [ tid + 2]; if (blocksize >= 2) s v[tid] 4-= sv [ tid + i]; } Spuštění kernelu: reduce6»(d_v ) ; Přenos 50.64 GB/s, 12.66 MEIem/s. Jiří Filipovič Optimalizace pro GPU hardware 49 / 52 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo ooooooooo oo ooooooooooooooo»oo Implementace 7 Můžeme algoritmus ještě vylepšit? Vratme se zpět ke složitosti: • celkem log n kroků • celkem n — 1 sčítání • časová složitost pro p threadů běžících paralelně (p procesorů) 0(* + \ogn) Cena paralelního výpočtu • definována jako počet procesorů krát časová složitost • přidělíme-li každému datovému elementu jeden thread, lze uvažovat p = n • pak je cena 0(n • log n) • není efektivní Jiří Filipovič Optimalizace pro GPU hardware 50 / 52 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo ooooooooo oo oooooooooooooooo«o Implementace 7 Snížení ceny • použijeme threadů • každý thread provede O(logn) sekvenčních kroků • následně se provede O(logn) paralelních kroků • časová složitost zůstane • cena se sníží na O(n) Co to znamená v praxi? • redukujeme práci spojenou s vytvářením threadu a pointerovou aritmetikou • to přináší výhodu v momentě, kdy máme výrazně více threadů, než je třeba k saturaci GPU • navíc snižujeme overhead spouštění kernelů Jiří Filipovič Optimalizace pro GPU hardware 51 /52 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO OOOOOOOOOOOOOOOOOO OOOOOOOOO OO 00000000000000000» Implementace 7 Modifikujeme načítání do sdílené paměti unsigned int gridSize = blockSize*2*gridDim.x; sv[tid] = 0; while(i < n){ sv[tid] += v[i] + v[i+blockSize ] ; i += gridSize; } __syncthreads (); Přenos 77.21 GB/s, 19.3MEIem/s. Jiří Filipovič Optimalizace pro GPU hardware 52 /52 Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO OOOOOOOOOOOOOOOOOO OOOOOOOOO OO 00000000000000000» Implementace 7 Modifikujeme načítání do sdílené paměti unsigned int gridSize = blockSize*2*gridDim.x; sv[tid] = 0; while(i < n){ sv[tid] += v[i] + v[i+blockSize ] ; i += gridSize; } __syncthreads (); Přenos 77.21 GB/s, 19.3MEIem/s. Jednotlivé implementace jsou k nalezení v CUDA SDK. Jiří Filipovič Optimalizace pro GPU hardware 52 /52