Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy oooo oooooooooooooooo ooooooooo oo oooooooooooooooooo ooooo Optimalizace pro GPU hardware J-w i- ■ i " "v in Fihpovic jaro 2017 Jiří Filipovič Optimalizace pro GPU hardware 1/55 Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy •ooo oooooooooooooooo ooooooooo oo oooooooooooooooooo ooooo 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í) J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy o«oo oooooooooooooooo ooooooooo oo oooooooooooooooooo ooooo 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 o 32 thredů (tzv. warp) musí provádět stejnou instrukci J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy o«oo oooooooooooooooo ooooooooo oo oooooooooooooooooo ooooo 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 o 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 J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy o«oo oooooooooooooooo ooooooooo oo oooooooooooooooooo ooooo 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 o 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 M u Iti procesor 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/55 Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy oo«o oooooooooooooooo ooooooooo oo oooooooooooooooooo ooooo Vlastnosti threadů Oproti CPU threadům jsou GPU thready velmi lehké", o 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 multi procesor je omezen. J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy ooo» oooooooooooooooo ooooooooo oo oooooooooooooooooo ooooo b GF 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 J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy ooo» oooooooooooooooo ooooooooo oo oooooooooooooooooo ooooo 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 J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy ooo» oooooooooooooooo ooooooooo oo oooooooooooooooooo ooooo 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ý o 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 J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy ooo» oooooooooooooooo ooooooooo oo oooooooooooooooooo ooooo a 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ý o 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/55 Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO »000000000000000 ooooooooo oo oooooooooooooooooo ooooo OOOO »000000000000000 ooooooooo oo Optimalizace přístupu do globální paměti Rychlost globální paměti se snadno stane bottleneckem • sirka 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 9 k paměti je nutno přistupovat sdruženě (coalescing) o je vhodné vyhnout se užívání pouze podmnožiny paměťových regionů (partition camping) J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO O0OOOOOOOOOOOOOO ooooooooo oo oooooooooooooooooo ooooo 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 • jedna paměťová transakce může přenášet 32-, 64-, nebo 128-bytová slova • polovina warpu (16 threadů) přistupuje do paměti paralelně • 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 J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy oooo oo«ooooooooooooo ooooooooo oo oooooooooooooooooo ooooo 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 nejnižším počtem transakcí • pořadí threadů může být vzhledem k přenášeným elementům libovolně permutované J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOO0OOOOOOOOOOOO ooooooooo oo oooooooooooooooooo ooooo 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. j in Fihpovic Optimalizace pro GPU hardware sVca\á 1.2) provádí dva přenosy. O 100 80 Edu < 40 GTX260 FX56O0 6 B 10 12 14 16 Offset J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy oooo ooooooo«oooooooo ooooooooo oo oooooooooooooooooo ooooo 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á. n i i i i i i i i i D 2 4 6 S 10 12 14 16 18 Stride Jiří Filipovič Optimalizace pro GPU hardware 13/55 Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOO0OOOOOOO ooooooooo oo oooooooooooooooooo ooooo Přístup do globální paměti u Fermi (c.c. > 2.0) a 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? o programy s nepředvídatelnou datovou lokalitou mohou běžet efektivněji • nezarovnaný přístup - v principu žádné zpomalení <* 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) Jiří Filipovič Optimalizace pro GPU hardware 14/55 Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy oooo ooooooooo«oooooo ooooooooo oo oooooooooooooooooo ooooo 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 9 implementován broadcast - pokud všichni přistupují ke stejnému údaji v paměti J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy oooo oooooooooo«ooooo ooooooooo oo oooooooooooooooooo ooooo Konflikty bank a 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 • zpomalení 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 16/55 Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOO0OOOO Přístup bez konfliktů / . / Thread 0 Bank 0 Thread 1 Bank 1 Thread 2 Bank 2 Thread 3 Bank 3 Th read 4 Bank 4 Thread 5 Bank 5 Th rea d 6 Bank 6 Thread 7 Bank 7 A Thread S Bank B Thread 3 Bank 9 Thread 10 Bank 10 Thread 11 Bank 11 Á Thread 12 Bank 12 Thread 1 3 Bank 13 Á Thread 14 Bank 14 Thread 1 5 Bank 15 / l □ rS1 = Jiří Filipovič Optimalizace pro GPU hardware 17/55 Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOO0OOO ooooooooo oo oooooooooooooooooo ooooo Vícecestné konflikty Jiří Filipovič Optimalizace pro GPU hardware 18/55 Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOO0OO ooooooooo oo oooooooooooooooooo ooooo Broadcast Jiří Filipovič Optimalizace pro GPU hardware 19/55 Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOO0O ooooooooo oo oooooooooooooooooo ooooo 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]; Jin Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO 000000000000000« ooooooooo oo oooooooooooooooooo ooooo Komunikace CPU-GPU Přenosy mezi systémovou a grafickou pamětí • je nutné je minimalizovat (často i za cenu neefektivní části výpočtu na GPU) o mohou být zrychleny pomcí page-locked paměti o je výhodné přenášet větší kusy současně • je výhodné překrýt výpočet s přenosem J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO «00000000 oo oooooooooooooooooo ooooo ranspozice 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 x = blockldx.x * blockDim.x + threadldx.x; int y = blockldx.y * blockDim.y + threadldx.y; odata[x*n + y] = idata[y*n + x]; } int n){ Jiří Filipovič □ [31 Optimalizace pro GPU hardware 22/55 Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO O0OOOOOOO oo oooooooooooooooooo ooooo 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 23/55 Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO O0OOOOOOO oo oooooooooooooooooo ooooo b GF 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. J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy oooo oooooooooooooooo oo«oooooo oo oooooooooooooooooo ooooo Odstranění prokládání Matici můžeme zpracovávat po dlaždicích 9 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 o díky tomu je jak čtení, tak zápis bez prokládání J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy oooo oooooooooooooooo oo«oooooo oo oooooooooooooooooo ooooo Odstranění prokládání Matici můžeme zpracovávat po dlaždicích 9 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 o 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 o 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) o nejvhodnější velikost určíme experimentálně Jiří Filipovič Optimalizace pro GPU hardware 24/55 Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO 000*00000 oo oooooooooooooooooo ooooo 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 += BLOCK_ROWS) tile[threadldx.y+i][threadldx.x] = idata[index_in+i*n] ; __syncthreads () ; for (int i = 0; i < TILE_DIM; i += BLOCK_ROWS) odata[index_out+i*n] = tile[threadldx.x][threadldx.y+i]; J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO OOOO0OOOO oo oooooooooooooooooo ooooo b GF 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. J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy oooo oooooooooooooooo oooo«oooo oo oooooooooooooooooo ooooo GF 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í □ Jiří Filipovič Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO OOOO0OOOO oo oooooooooooooooooo ooooo b GF 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 J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO OOOO0OOOO oo oooooooooooooooooo ooooo b GF 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í J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO OOOOO0OOO oo oooooooooooooooooo ooooo b GF 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 ] ; J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO OOOOO0OOO oo oooooooooooooooooo ooooo 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. J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO OOOOO0OOO oo oooooooooooooooooo ooooo 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. Řešením je padding: __shared__ float tile[TILE_DIM][TILE_DIM + 1]; J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO OOOOOO0OO oo oooooooooooooooooo ooooo b GF Nyní dosahuje naše implementace výkon 93.4 GB/s. 9 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) J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO OOOOOOO0O oo oooooooooooooooooo ooooo 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 můžeme ztratit výhodu vyššího výkonu GPU... J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO 00000000« oo oooooooooooooooooo ooooo 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í J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO OOOOOOOOO 90 oooooooooooooooooo ooooo Rychlost provádění instrukcí Některé instrukce jsou v porovnání s ostatními pomalejší, než u procesoru 9 celočíselné delenia modulo • 32-bitové násobení celých čísel u c.c. 1.x • 24-bitové násobení celých čísel 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. J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy oooo oooooooooooooooo ooooooooo om oooooooooooooooooo ooooo Smyčky Malé cykly mají značný overhead • je třeba provádět podmíněné skoky o 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 u n roli J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO OOOOOOOOO OO »00000000000000000 ooooo b GF Součet prvků vektoru Pro vektor v o n prvcích chceme spočítat x = X)ľ=i Vi J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO OOOOOOOOO OO «00000000000000000 ooooo oučet prvků vektoru Pro vektor v o n prvcích chceme spočítat x = Y11=i v> Zápis v jazyce C int x = 0; for (int i = 0; i < n; i++) x += v [ i ] ; Jednotlivé iterace cyklu jsou na sobě závislé. Jiří Filipovič S1 Optimalizace pro GPU hardware "O ^ O' Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO OOOOOOOOO OO »00000000000000000 ooooo Součet prvků vektoru Pro vektor v o n prvcích chceme spočítat x = yí-Zápis v jazyce C int x = 0; 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í o není tedy nutno počítat sekvenčně J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy oooo oooooooooooooooo ooooooooo oo o«oooooooooooooooo ooooo b GF Paralelní algoritmus Představený sekvenční algoritmus provádí pro 8 prvků výpočet: (((((("i + v2) + "ä) + v4) + v$) + W>) + vr) + v8 j in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy oooo oooooooooooooooo ooooooooo oo o«oooooooooooooooo ooooo Paralelní algoritmus Představený sekvenční algoritmus provádí pro 8 prvků výpočet: (((((("i + v2) + "ä) + v4) + v$) + W>) + vr) + v8 Sčítání je asociativní... spřeházejme tedy závorky: ((^1 + v2) + (vs + v4)) + ((1/5 + v6) + (v7 + v8)) Jin Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy oooo oooooooooooooooo ooooooooo oo o«oooooooooooooooo ooooo a Paralelní algoritmus Představený sekvenční algoritmus provádí pro 8 prvků výpočet: ((((((ví + V2) + v3) + v4) + vb) + v6) + vj) + v8 Sčítání je asociativní... spřeházejme tedy závorky: ((^1 + v2) + (v3 + v4)) + ((vb + v6) + (w + v8)) Nyní můžeme pracovat paralelně o v prvním kroku provedeme 4 sčítání 9 ve druhém dvě 9 ve třetím jedno Celkově stejné množství práce (r? — 1 sčítání), ale v log2 n paralelních krocích! Jiří Filipovič Optimalizace pro GPU hardware 34/55 Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO OOOOOOOOO OO OO0OOOOOOOOOOOOOOO ooooo Paralelní algoritmus a 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ů o vyžaduje globální bariéru J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO OOOOOOOOO OO OOO0OOOOOOOOOOOOOO ooooo b GF Naivní přístup Nejjednodušší schéma algoritmu (r?je mocnina dvou): • kernel pro sudá / < n provede v[i] += v[i+l] o 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. J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO OOOOOOOOO OO 0000*0000000000000 ooooo b GF Využití rychlejší paměti V rámci volání kernelu můžeme posčítat více, než jen dvojice 9 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) o 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ů m mz rn'°Smn V / m—l • približne n + — ctem, — zápisu r m ' m r 9 \ogm n spuštění kernelu j in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy oooo oooooooooooooooo ooooooooo oo ooooo«oooooooooooo ooooo mplementace __global__ void reducel(int *v){ extern __shared__ int sv [ ]; unsigned int tid = threadldx.x; unsigned int i = blockldx.x*blockDim.x + threadldx.x; sv[tid] = 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]; Jin Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO OOOOOOOOO OO OOOOOO0OOOOOOOOOOO ooooo Vysoká úroveň divergence 9 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.77 GB/s, 0.94 MElem/s. □ s Jiří Filipovič Optimalizace pro GPU hardware 39/55 Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO OOOOOOOOO OO OOOOOOO0OOOOOOOOOO ooooo Implementace 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.33 GB/s, 2.08 MEIem/s. ■v Reší divergenci, generuje konflikty bank. s * = 2) { Jiří Filipovič Optimalizace pro GPU hardware 5 ^(^(V 40/55 Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO OOOOOOOOO OO OOOOOOOO0OOOOOOOOO ooooo Implementace 3 Tak ještě jinak... for (unsigned int s = blockDim . x/2; s > 0; s »= if (tid < s) s v [ t i d ] += s v [ t i d + s ] ; __syncthreads(); } Žádná divergence ani konflikty. Přenos 16.34 GB/s, 4.08 MElem/s. Polovina threadů nic nepočítá... i) { Jiří Filipovič Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO OOOOOOOOO OO 000000000*00000000 ooooo Implementace a 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. J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO OOOOOOOOO OO OOOOOOOOOO0OOOOOOO ooooo Implementace a 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 9 musíme sčítat přes volatile proměnnou o podmínka if(tid < s) je zde zbytečná (nic neušetří) Můžeme unrollovat poslední warp. Pozor, jedná se o porušení obecné korektnosti kódu (funkční pouze na HW s velikostí warpu dělitelnou 32, což jsou zatím všechna NVIDIA i AMD GPU). Jiří Filipovič Optimalizace pro GPU hardware 43/55 Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy oooo oooooooooooooooo ooooooooo oo ooooooooooo«oooooo ooooo Implementace 5 float mySum = 0; for (unsigned int s = blockDim . x / 2; s > 32; s »= 1){ if (tid < s) sv[tid] = mySum = mySum + sv[tid + s]; __syncthreads(); } if (tid < 32){ volat ile float * s = SV > s tid] mySum = mySum + s tid + 32] s tid] mySum = mySum + s tid + 16] s tid] mySum = mySum + s tid + 8]; s tid] mySum = mySum + s tid + 4]; s tid] mySum = mySum + s tid + 2]; s tid] mySum = mySum + s tid + i]; } Ušetříme čas i ostatním warpům (skončí dříve s for cyklem). Přenos 37.68 GB/s, 9.42 MEIem/s. Jiří Filipovič Optimalizace pro GPU hardware 44/55 Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO OOOOOOOOO OO OOOOOOOOOOOO0OOOOO ooooo mplementace 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í? o algoritmus pracuje s bloky o velikosti 2n • velikost bloku je shora omezena o známe-li při kompilaci velikost bloku, můžeme použít šablonu template __global__ void reduce6(int *v) J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO OOOOOOOOO OO OOOOOOOOOOOOO0OOOO ooooo a mplementace Podmínky s blockSize se vyhodnotí již pří překladu if (blockSize >= 512){ if (tid < 256) sv[tid] = mySum = mySum + sv[tid + 256]; __syncthreads(); } if (blockSize >= 256){ if (tid < 128) sv[tid] = mySum = mySum + sv[tid + 128]; __syncthreads(); } if (blockSize >= 128){ if (tid < 64) sv[tid] = mySum = mySum + sv[tid + 64]; __syncthreads(); } J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy oooo oooooooooooooooo ooooooooo oo oooooooooooooo«ooo ooooo mplementace 6 Spuštění kernelu: reduce6(d_v ) ; Přenos 50.64GB/s, 12.66 MElem/s. Jiří Filipovič S1 Optimalizace pro GPU hardware "O ^ O' Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO OOOOOOOOO OO OOOOOOOOOOOOOOO0OO ooooo a Implementace 7 Můžeme algoritmus ještě vylepšit? Vraťme se zpět ke složitosti: 9 celkem log n kroků • celkem n — 1 sčítání • časová složitost pro p threadů běžících paralelně (p procesorů) Cena paralelního výpočtu 9 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) 9 není efektivní Jiří Filipovič Optimalizace pro GPU hardware 48/55 Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO OOOOOOOOO OO OOOOOOOOOOOOOOOO0O ooooo a Implementace 7 Snížení ceny • použijeme 0(^^) threadů • každý thread provede 0(\ogn) sekvenčních kroků • následně se provede 0(\ogn) 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 9 to přináší výhodu v momentě, kdy máme výrazně více threadů, než je třeba k saturaci GPU o navíc snižujeme overhead spouštění kernelů Jiří Filipovič Optimalizace pro GPU hardware 49/55 Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce oooo oooooooooooooooo ooooooooo oo ooooooooooooooooot mplementace Modifikujeme načítání do sdílené paměti unsigned int gridSize = blockSize*2*gridDim . x ; sv[tid] = 0; while(i < n){ sv[t i d] += v[i] + v[i+blockSize]; i += gridSize; } __syncthreads () ; Přenos 77.21 GB/s, 19.3 MElem/s. Jiří Filipovič S1 Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce oooo oooooooooooooooo ooooooooo oo ooooooooooooooooo< mplementace Modifikujeme načítání do sdílené paměti unsigned int gridSize = blockSize*2*gridDim.x ; sv[tid] = 0; while(i < n){ sv[t i d] += 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 J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO OOOOOOOOO OO OOOOOOOOOOOOOOOOOO »0000 GPU vs. CPU GPU vhodnější pro výpočty omezené aritmetickou či paměťovou propustností • těží z vyšších teoretických maxim (pozor na sběrnici u výpočtů omezených paměťovou propustností) o násobení matic • metoda konečných diferencí CPU vhodnější pro výpočty omezené latencí • komplexní čip je více flexibilní, více výkonu na vlákno (pro nedostatečně paralelizovatelné problémy), rychlejší náhodný přístup do paměti • procházení grafu Jiří Filipovič Optimalizace pro GPU hardware 51 / 55 Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy oooo oooooooooooooooo ooooooooo oo oooooooooooooooooo o«ooo Paralelizace Sčítání vektorů • jednoduché datově-paralelní vyjádření • žádná synchronizace o potřebujeme velké vektory Redukce • na první pohled může vypadat sekvenčně a ve skutečnosti realizovatelná v logr? krocích o často je třeba nedržet se sekvenční verze a zamyslet se nad paralelizací problému (ne sekvenčního algoritmu) j in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy OOOO OOOOOOOOOOOOOOOO OOOOOOOOO OO OOOOOOOOOOOOOOOOOO OO0OO Divergence kódu Divergence kódu • serializace, divergují-li thready uvnitř warpu o nalezení nedivergujícího algoritmu může být snadné • redukce o ale také může prakticky znemožnit akceleraci některých jinak dobře paralelizovatelných algoritmů 9 mnoho nezávislých stavových automatů • nutnost zamyslet se nad výrazně odlišným algoritmem pro daný problém Jiří Filipovič Optimalizace pro GPU hardware 53/55 Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy oooo oooooooooooooooo ooooooooo oo oooooooooooooooooo oooto Divergence přístupu do paměti Divergence přístupu do paměti 9 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 9 může vyžadovat využití odlišných datových struktur 9 práce s řídkými maticemi 9 u rigidnějších struktur si lze často pomoci on-chip pamětí • transpozice matic J in Fihpovic Optimalizace pro GPU hardware Paralelismus Optimalizace přístupu do paměti Transpozice matic Skalární optimalizace Redukce Algoritmy oooo oooooooooooooooo ooooooooo oo oooooooooooooooooo oooo* Latence GPU GPU je dnes často propojena se zbytkem systému přes PCI-E 9 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 J in Fihpovic Optimalizace pro GPU hardware