Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo ooooooooo oo oooooooooooooooooo Výkon GPU hardware Jiří Filipovič podzim 2010 Jiří Filipovič Výkon 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í) □ - = = ^Q^O Jiří Filipouič Výkon GPU hardware 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 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 □ - = = ^Q^O Jiří Filipouič Výkon GPU hardware 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 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č Výkon GPU hardware 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. □ - = = ^Q^O Jiří Filipouič Výkon GPU hardware 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 Výkon GPU hardware 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 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 □ - = = ^Q^O Jiří Filipouič Výkon GPU hardware 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. □ - = = ^Q^O Jiří Filipovič Výkon GPU hardware 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á (G200 > 24 flops/float, G100 > 30) • 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 spojitě (coalescing) • je vhodné vyhnout se užívání pouze podmnožiny paměťových regionů (partition camping) Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOO Transpozice matic ooooooooo Rychlost instrukcí 00 Spojitý 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 > 128B aligned segment ......I.......... Half warp of threads Výkon GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oo»ooooooooooooooo ooooooooo oo oooooooooooooooooo Spojitý 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 □ - = = ^Q^O Jiří Filipouič Výkon GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo ooo«oooooooooooooo ooooooooo oo oooooooooooooooooo Spojitý 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 nejnižším počtem transakcí • pořadí threadů může být vzhledem k přenášeným elementům libovolně permutované Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooo»ooooooooooooo ooooooooo oo oooooooooooooooooo Spojitý přístup do paměti (c.c. < 2.0) Thready jsou zarovnané, blok elementů souvislý, pořadí není permutované - spojitý přístup na všech GPU. k ál 1 11 1 1 1 1 1 1 1 1 1 1 1 1 Výkon 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) Thready nejsou zarovnané, blok elementů souvislý, pořadí není permutované - jedna transakce na GPU s c.c. > 1.2. iVi'i'i'i'i'i'i'i'i'iŕŤWn Výkon GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce QOOO oooooo»ooooooooooo ooooooooo oo oooooooooooooooooo Nezarovnaný prístup do paměti(c.c. < 2.0) Obdobný případ může vézt k nutnosti použít dvě transakce. □ □ s - - Jiří Filipovie Výkon GPU hardware 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. 140 0 2 4 S 6 10 12 14 16 Offset Jiří Filipovič Výkon GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo 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-1—li i—i 0 2 4 B 8 10 12 14 16 18 Stride □ - = = ^Q^O Jiří Filipouič Výkon GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo ooooooooosoooooooo 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 It i 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 oooooooooo»ooooooo 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ří Filipouič Výkon GPU hardware 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 □ - = = ^Q^O Jiří Filipouič Výkon GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooo»ooooo ooooooooo oo oooooooooooooooooo Prístup bez konfliktů Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo ooooooooooooo»oooo ooooooooo oo oooooooooooooooooo Vřcecestné konflikty Jiří Filipovič Výkon GPU hardware Paralelismus Optimalizace přístupu do paměti oooo oooooooooooooooooo Transpozice matic ooooooooo Rychlost instrukcí 00 Redukce oooooooooooooooooo Jiří Filipovič Výkon GPU hardware 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]; Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooo»o 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í zarovnaného přístupu • nevhodná, pokud je bottleneck latence • může zjednodušit adresování či přidat filtraci Jiří Filipovič Výkon GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo ooooooooooooooooo* ooooooooo oo oooooooooooooooooo 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 □ - = = ^Q^O Jiří Filipouič Výkon GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO OOOOOOOOOOOOOOOOOO »00000000 oo oooooooooooooooooo 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]; 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? 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. 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 blocích • načteme po řádcích blok do sdílené paměti • uložíme do globální paměti jeho transpozici taktéž po řádcích • díky tomu je jak čtení, tak zápis bez prokládání 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 blocích • načteme po řádcích blok do sdílené paměti • uložíme do globální paměti jeho transpozici taktéž po řádcích • díky tomu je jak čtení, tak zápis bez prokládání Jak velké bloky použít? • budeme uvažovat bloky čtvercové velikosti • pro zarovnané čtení musí mít řádek bloku velikost dělitelnou 16 • v úvahu připadají bloky 16 x 16, 32 x 32 a 48 x 48 (jsme omezeni velikostí sdílené paměti) • nejvhodnější velikost určíme experimentálně □ - = = ^Q^O Jiří Filipovič Výkon GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo ooo»ooooo oo oooooooooooooooooo Bloková 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 += BLOCK.ROMS) tile[threadldx.y+i][threadldx.x] = idata[index_in+i*n]; __syncthreads(); for (int i = 0; i < TILE.DIM; i += BLOCK.ROMS) odata[index_out+i*n] = tile[threadldx.x][threadldx. y+i] ; } □ SP - = .g -o<^o Jiří Filipouič Výkon GPU hardware Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOO Transpozice matic oooo»oooo Rychlost instrukcí 00 Nejvyšší výkon byl naměřen při použití bloků velikosti 32 velikost thread bloku 32 x 8, a to 75.1GB/s. 32, Výkon GPU hardware Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOO Transpozice matic oooo»oooo Rychlost instrukcí 00 Nejvyšší výkon byl naměřen při použití bloků 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í s Výkon GPU hardware Optimalizace přístupu do paměti oooooooooooooooooo Transpozice matic oooo»oooo Rychlost instrukcí 00 Nejvyšší výkon byl naměřen při použití bloků 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 s Výkon GPU hardware Optimalizace přístupu do paměti oooooooooooooooooo Transpozice matic oooo»oooo Rychlost instrukcí 00 Nejvyšší výkon byl naměřen při použití bloků 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 v rámci bloků pouze kopírujeme, dosáhneme výkonu 94.9GB/s • něco ještě není optimální s Výkon GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo ooooo»ooo oo oooooooooooooooooo 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] ; Jiří Filipouič Výkon GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo ooooo»ooo oo oooooooooooooooooo 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. Jiří Filipouič Výkon GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo ooooo»ooo oo oooooooooooooooooo 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]; 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) Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce OOOO OOOOOOOOOOOOOOOOOO 0000000*0 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... Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo oooooooo* 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í □ - = = ^Q^O Jiří Filipouič Výkon GPU hardware 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), __expr7xj, sincosf(x), ^rsqrtf(x) aj. □ - = = ^Q^O Jiří Filipouič Výkon GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo ooooooooo o» 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 □ - = = ^Q^O Jiří Filipouič Výkon 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 = ^"=1 v-,. Jiří Filipouič Výkon 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 = J^/Li 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é. Výkon 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 = J^/Li 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 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: (((((( ví + ^2) + v3) + 1/4) + v5) + v6) + vj) + v8 Výkon 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: (((((( ví + ^2) + v3) + 1/4) + v5) + v6) + vj) + v8 Sčítání je asociativní... spřeházejme tedy závorky: ((1/1 + v2) + (1/3 + vA)) + ((1/5 + v6) + {vj + v8)) Výkon 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: (((((( ví + v2) + v3) + 1/4) + v5) + v6) + vj) + v8 Sčítání je asociativní... spřeházejme tedy závorky: ((1/1 + v2) + (1/3 + vA)) + ((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č Výkon GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo ooooooooo oo oo»ooooooooooooooo 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 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á / < 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. □ g - = = ^o^O Jiří Filipouič Výkon GPU hardware 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í YlT^mXx vi Výhodnější z hlediska paměťových přenosů i spouštění kernelů • přibližně n + čtení, ^ zápisů • logm n spuštění kernelu 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[]; 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 *= 2) { if (tid % (2*s) == 0) sv[tid] += sv[tid + s]; __syncthreads(); } if (tid = 0) v[blockldx.x] = sv[0]; □ 0; s »= 1) { if (tid < s) sv[tid] += sv[tid + s]; __ sync t hr e ads (); } Žádná divergence ani konflikty. Přenos 16.34 GB/s, 4.08 MElem/s. Polovina threadů nic nepočítá... 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 = blockldx.x*(blockDim.x*2) + threadldx.x; sv[tid] = v[i] + v[i+blockDim.x]; Přenos 27.16 GB/s, 6.79 MElem/s. Data zřejmě čteme optimálně, stále je zde však výkonová rezerva -zaměřme se na instrukce. 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... 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) sv[tid] += sv[tid + s]; __syncthreads (); } if (tid < 32){ sv tid] += sv tid ^ - 32] sv tid] += sv tid ^ - 16] sv tid] += sv tid ^ - 8]; sv tid] += sv tid ^ " 4]; sv tid] += sv tid ^ " 2]; sv tid] += sv tid ^ - i]; } Ušetříme čas i ostatním waprům (zkončí dříve s for cyklem). Přenos 37.68 GB/s, 9.42 MElem/s. Jiří Filipovič Výkon GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo ooooooooo oo oooooooooooosooooo 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) □ - = = ^Q^O Jiří Filipouič Výkon GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí Redukce oooo oooooooooooooooooo ooooooooo oo ooooooooooooo»oooo Implementace 6 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) sv[tid] += sv[tid + 128]; __syncthreads (); } if (blockSize >= 128){ if (tid < 64) sv[tid] += sv[tid + 64]; __syncthreads (); □ - = = -0*3*0 Jiří Filipouič Výkon GPU hardware 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) sv[t id] += sv[t id 4- 32] if (blockSize >= 32) sv[t id] += sv[t id 4- 16] if (blockSize >= 16) sv[t id] += sv[t id 4- 8]; if (blockSize >= 8) sv[t id] ■f= sv [ tid + 4]; if (blockSize >= 4) sv[t id] ■f= sv[tid + 2]; if (blockSize >= 2) sv[t id] ■f= sv [ tid + l]: } Spustení kernelu: reduce6»(d_v ); Přenos 50.64GB/s, 12.66MEIem/s. Jiří Filipouič Výkon GPU hardware 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 logn kroků • celkem n — 1 sčítání • časová složitost pro p threadů běžících paralelně (p procesorů) 0(£ + logn) 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 rí) • není efektivní □ g - = = ^q^O Jiří Filipovič Výkon GPU hardware 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 O(j^) threadů • každý thread provede (D(logn) sekvenčních kroků • následně se provede (D(logn) paralelních kroků • časová složitost zůstane • cena se sníží na 0{rí) 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ů □ - = = ^Q^O Jiří Filipovič Výkon GPU hardware 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ří Filipouič Výkon GPU hardware 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.3 MElem/s. Jednotlivé implementace jsou k nalezení v CUDA SDK. □ S ~ = -š -00,0 Jiří Filipouič Výkon GPU hardware