Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí OOOOOOOOOOOOOOOOOOOOO OOOOOOOOOOOOOOOOO OOOOOOOOO Výkon GPU hardware Jiří Filipovič podzim 2014 Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti ♦OOOOOOOOOOOOOOOOOOOO Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOOO 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, GK110 > 62, GM204 > 82) • 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) • je vhodné vyhnout se užívání pouze podmnožiny paměťových regionů (partition camping) Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti o«ooooooooooooooooooo Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOOO 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 Optimalizace přístupu do paměti 00*000000000000000000 Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOOO 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č Výkon GPU hardware Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí OOO0OOOOOOOOOOOOOOOOO OOOOOOOOOOOOOOOOO OOOOOOOOO 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 přístupu do paměti Transpozice matic Rychlost instrukcí OOOO^OOOOOOOOOOOOOOOO OOOOOOOOOOOOOOOOO OOOOOOOOO Sdružený přístup do paměti (c.c. < 2.0) Thready jsou zarovnané, blok elementů souvislý, pořadí není permutované - efektivní přístup na všech GPU. á i J 1 1 I 1 1 1 1 r t 1 1 1 1 Optimalizace přístupu do paměti OOOOO0OOOOOOOOOOOOOOO Transpozice matic OOOOOOOOOOOOOOOOO Nezarovnaný přístup do paměti (c.c. < 2.0) Rychlost instrukcí OOOOOOOOO Thready nejsou zarovnané, blok elementů souvislý, pořadí není permutované - jedna transakce na GPU s c.c. > 1.2. 'ľľľľľľl Výkon GPU hardware Optimalizace přístupu do paměti oooooo«oooooooooooooo Transpozice matic OOOOOOOOOOOOOOOOO Nezarovnaný přístup do paměti (c.c. < 2.0) Rychlost instrukcí OOOOOOOOO Optimalizace přístupu do paměti ooooooo»ooooooooooooo Transpozice matic ooooooooooooooooo Rychlost instrukcí ooooooooo 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. GPU (c.c. > 1.2) provádí dva přenosy. -W-GTX280 ? -*FX5600 2 4 6 8 10 12 14 16 Offset i -00.0 Výkon GPU hardware Optimalizace přístupu do paměti OOOOOOOOÄOOOOOOOOOOOO Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOOO 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č Výkon GPU hardware Optimalizace přístupu do paměti OOOOOOOOO^OOOOOOOOOOO Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOOO Přístup do globální paměti u Fermi (c.c. > 2.0) Fermi má LI a L2 cache • LI: 128 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) Optimalizace přístupu do paměti OOOOOOOOOOÄOOOOOOOOOO Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOOO Přístup do globální paměti u herní varianty Kepler (c.c. = 3.0) „Herní" Kepler používá pro obecný přístup pouze L2 cache • LI: pouze pro lokální paměť, celkem 16 KB, 32 KB nebo 48 KB • L2: 32 byte na řádek, až 1.5 GB na GPU Optimalizace přístupu do paměti OOOOOOOOOOO^OOOOOOOOO Rychlost instrukcí OOOOOOOOO Transpozice matic OOOOOOOOOOOOOOOOO Přístup do globální paměti u Kepler (c.c. = 3.5) a Maxwe „Plnotučný" Kepler a Maxwell má vedle L2 cache ještě Data cache • sdílená s texturami, podporuje c.c. > 3.5 • pro data pouze ke čtení • může rozeznat kompilátor, pomůžeme mu pomocí const __restrict__ či explicitně __ldg() Maxwell nepoužívá LI cache pro lokální paměť • kódy které ji hojně využívají je třeba reoptimalizovat Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti OOOOOOOOOOOOÄOOOOOOOO Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOOO Partition camping • relevantní pro c.c. 1.x • procesory založené na G80 mají 6 regionů, G200 mají 8 regionů globální paměti • paměť je dělena do regionů po 256-bytech • pro maximální výkon je zapotřebí, aby bylo přistupováno rovnoměrně k jednotlivým regionům • mezi jednotlivými bloky • ty se zpravidla spouští v uspořádání daném polohou bloku v gridu • pokud je využívána jen část regionů, nazýváme jev partition camping • obecně ne tak kritické, jako spojitý přístup • záludnější, závislé na velikosti problému Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti ooooooooooooo«ooooooo Transpozice matic ooooooooooooooooo HW organizace sdílené paměti Rychlost instrukcí OOOOOOOOO 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.0 32 bank, paměťový prostor mapován prokládané s odstupem 32 bitů či 64 bitů (c.c. 3.x) • 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č Výkon GPU hardware Optimalizace přístupu do paměti OOOOOOOOOOOOOOÍOOOOOO Transpozice matic ooooooooooooooooo Rychlost instrukcí ooooooooo 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 několika výjimkami • thready přistupují ke stejnému místu v paměti o thready přistupují k rozdílným podslovům 64-bitového slova (c.c. > 3.0) • 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 i -00.0 Výkon GPU hardware Optimalizace přístupu do paměti ooooooooooooooo«ooooo Transpozice matic ooooooooooooooooo Rychlost instrukcí ooooooooo Přístup bez konfliktů Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOÄOOOO Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOOO Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti ooooooooooooooooo«ooo Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOOO Broadcast Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOO0OO Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOOO Vzory přístupu Zarovnání není třeba, negeneruje konflikty bank int x = s[threadldx.x + offset]; Prokládání negeneruje konflikty, je-li c liché, u c.c. > 3.0 může být c = 2 (ne pro 64-bitová čísla) int x = s[threadldx.x * c]; Přístup ke stejné proměnné negeneruje na c.c. > 2.0 konflikty nikdy, na 1,x je-li počet c threadů přistupující k proměnné násobek 16 int x = s[threadldx.x / c]; Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOOOÄO Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOOO 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č Výkon GPU hardware Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOOOO* Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOOO 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 přístupu do paměti ooooooooooooooooooooo Transpozice matic ♦OOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOOO 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]; } Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti ooooooooooooooooooooo Výkon Transpozice matic OÄOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOOO 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? i -00.0 Výkon GPU hardware Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí ooooooooooooooooooooo o«ooooooooooooooo ooooooooo 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. Optimalizace přístupu do paměti ooooooooooooooooooooo Transpozice matic oo»oooooooooooooo Rychlost instrukcí ooooooooo 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í Optimalizace přístupu do paměti ooooooooooooooooooooo Transpozice matic oo»oooooooooooooo Rychlost instrukcí ooooooooo 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 sdružené č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ě Optimalizace přístupu do paměti ooooooooooooooooooooo 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 += 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]; } Transpozice matic OOOÄOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOOO Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOOOOO Transpozice matic OOOO^OOOOOOOOOOOO Rychlost instrukcí OOOOOOOOO 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. i -00.0 Výkon GPU hardware Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOOOOO Transpozice matic OOOO^OOOOOOOOOOOO Rychlost instrukcí OOOOOOOOO Nej vyšší 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í i -00.0 Výkon GPU hardware Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOOOOO Transpozice matic OOOO^OOOOOOOOOOOO Rychlost instrukcí OOOOOOOOO Nej vyšší 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 N ej vyšší 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í Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOOOOO Sdílená paměť Transpozice matic OOOOOÄOOOOOOOOOOO Rychlost instrukcí OOOOOOOOO 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]; m -00,0 Výkon GPU hardware Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOOOOO Sdílená paměť Transpozice matic OOOOOÄOOOOOOOOOOO Rychlost instrukcí OOOOOOOOO 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. Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí OOOOOOOOOOOOOOOOOOOOO ooooo«ooooooooooo OOOOOOOOO 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 přístupu do paměti ooooooooooooooooooooo Výkon Transpozice matic 000000*0000000000 Rychlost instrukcí OOOOOOOOO 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 (viz. partition camping) Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOOOOO Transpozice matic OOOOOOOÄOOOOOOOOO Rychlost instrukcí OOOOOOOOO Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí ooooooooooooooooooooo oooooooo«oooooooo ooooooooo Výkon 500 1000 1500 2000 2500 3000 3500 4000 velikost matice Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí OOOOOOOOOOOOOOOOOOOOO OOOOOOOOOÄOOOOOOO OOOOOOOOO Poklesy výkonu Pro některé velikosti problému výkon klesá, v tomto chování lze nalézt pravidla Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti ooooooooooooooooooooo Poklesy výkonu Transpozice matic OOOOOOOOOÄOOOOOOO Rychlost instrukcí OOOOOOOOO Pro některé velikosti problému výkon klesá, v tomto chování lze nalézt pravidla • u matic o velikosti dělitelné 512 dosahujeme pouze cca 19GB/s • pro zbývající o velikosti dělitelné 256 cca 35 GB/s • pro zbývající o velikosti dělitelné 128 cca 62 GB/s Optimalizace přístupu do paměti ooooooooooooooooooooo Transpozice matic oooooooooo«oooooo Rychlost instrukcí ooooooooo Poklesy výkonu Jeden region paměti má šířku 2 bloků (256 byte / 4 byte na float, 32 floatů v bloku). Podíváme-li se na umístění bloků vzhledem k velikosti matice, zjistíme, že • při velikosti dělitelné 512 jsou bloky ve sloupcích ve stejném regionu • při velikosti dělitelé 256 je každý sloupec nejvýše ve dvou regionech • při velikosti dělitelné 128 je každý sloupec nejvýše ve čtyřech regionech Dochází tedy k partition campingu! Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí ooooooooooooooooooooo ooooooooooo»ooooo ooooooooo Jak odstraníme partition camping Můžeme doplnit „slepá data" a vyhnout se tak nevhodným velikostem matic • to komplikuje práci s algoritmem • další nevýhodou jsou větší paměťové nároky Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí ooooooooooooooooooooo ooooooooooo»ooooo ooooooooo Jak odstraníme partition camping Můžeme doplnit „slepá data" a vyhnout se tak nevhodným velikostem matic • to komplikuje práci s algoritmem • další nevýhodou jsou větší paměťové nároky Můžeme změnit mapování id thread bloků na bloky v matici • diagonální mapování zajistí přístup do rozdílných regionů int blockIdx_y = blockldx.x; int block!dx_x = (blockldx.x+blockldx.y) % gridDim.x; Nová implementace podává výkon cca 80 GB/s • ten neklesá na datech, kde klesal výkon původní implementace • pro matice o velikosti nedělitelné 128 je však nižší • algoritmus je složitější • můžeme jej však používat jen u dat, pro které je původní implementace nevýhodná Pro daný problém nemusí existovat (a spravidla neexistuje) ideální algoritmus pro celý rozsah (či typ) vstupních dat, je vhodné řádně benchmarkovat (ne každý potenciální problém odhalíme pohledem na kód). Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOOOOO Transpozice matic OOOOOOOOOOOOO0OOO Rychlost instrukcí OOOOOOOOO Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOOOOO Transpozice matic OOOOOOOOOOOOOOÄOO Rychlost instrukcí OOOOOOOOO Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí ooooooooooooooooooooo ooooooooooooooo»o ooooooooo 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... Optimalizace přístupu do paměti ooooooooooooooooooooo Transpozice matic OOOOOOOOOOOOOOOO* Rychlost instrukcí ooooooooo 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í memory 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č Výkon GPU hardware Optimalizace přístupu do paměti ooooooooooooooooooooo Provádění instrukcí Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí •OOOOOOOO Provádění instrukcí na m u Iti procesoru (c.c. 1.x) • zde je 8 SP jader a 2 SFU jádra • nedojde-li k překryvu SP a SFU provádění instrukcí, může multiprocesor dokončit až 8 instrukcí na takt • jeden warp je tedy proveden za 4 nebo více taktů • některé instrukce jsou výrazně pomalejší • znalost doby provádění instrukcí nám pomůže psát efektivní kód Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOOOOO Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí OÄOOOOOOO Operace v pohyblivé řádové čárce GPU je primárně grafický HW • v grafických operacích pracujeme zpravidla s čísly s plovoucí řádovou čárkou • GPU je schopno provádět je velmi rychle • novější GPU (compute capability > 1.3) dokážou pracovat i v double-precision, starší pouze v single-precision • některé aritmetické funkce jsou používány v grafických výpočtech velmi často • GPU je implementuje v hardware • HW implementace poskytuje méně přesné výsledky (pro spoustu aplikací není problém) • rozlišeno pomocí prefixu Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti ooooooooooooooooooooo Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí 00*000000 Aritmetické operace Operace s plovoucí řádovou čárkou • sčítání, násobení velmi rychlé • násobení a sčítání může být u c.c. 1.x kombinováno do jedné instrukce MAD o nižší přesnost • rychlost 1 cyklus na SP • _Jaddjrn() a __fmuLrn() lze použít pokud nechceme, aby byla v překladu použita instrukce MAD • MAD je nahrazeno FMAD u c.c. 2.x (shodná rychlost, vyšší přesnost) • 64-bitové verze pomalejší 1/8 (1.3), 1/2 (2.0), 1/12 (2.1), 1/24 (3.0), 1/3 (3.5), 1/32 (5.0) • dělení relativně pomalé, rychlejší je převrácená hodnota • rychlejší varianta pomocí —fdividef(x, y) • rychlá reciproká druhá odmocnina Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti ooooooooooooooooooooo Transpozice matic ooooooooooooooooo Rychlost instrukcí 000*00000 Aritmetické operace Operace s plovoucí řádovou čárkou • sinf(x), __cosf(x), __expf(x) • přesnější sinf(x), cosf(x), expf(x) řádově pomalejší • operací s různými rychlostmi a přesností je implementováno více, viz CUDA manuál Celočíselné operace • sčítání jako u plovoucí řádové čárky • násobení u c.c. 1.x pomalé • __mu/24(x, y) a __umul24(x, y) rychlejší • násobení u c.c. > 2.0, 24-bitová verze naopak pomalá • dělení a modulo velmi pomalé, pokud je n mocnina 2, můžeme využít • i/n je ekvivalentní / >> log2(n) o i%n je ekvivalentní i&i(n — 1) Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOOOOO Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí 0000*0000 Smyčky Malé cykly mají značný overhead • je třeba provádět skoky • je třeba vyhodnocovat podmínky • 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 Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOOOOO Ostatní operace Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOÄOOO Další běžné instrukce jsou prováděny základní rychlostí (tj. odpovídající počtu SP) • porovnávání • základní bitové operace (ne posuvy) • instrukce přistupující do paměti (s omezeními popsanými výše a s omezením latence a šířky pásma) • jako ofFset mohou použít hodnotu v registu + konstantu • synchronizace (pokud ovšem nečekáme :-)) Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti ooooooooooooooooooooo Transpozice matic ooooooooooooooooo Rychlost instrukcí 000000*00 Pozor na sdílenou paměť Pokud nedojde ke konfliktům bank, je sdílená paměť u c.c. 1.x rychlá jako registry. Ale pozor • instrukce dokáží pracovat pouze s jedním operandem ve sdílené paměťi • použijeme-li v rámci jedné instrukce více operandů ve sdílené paměti, je třeba explicitní load/store • instrukce MAD běží pomaleji (c.c. 1.x) • a + s[i] 4 cykly na warp • a + a * s[i] 5 cyklů na warp • a + b * s[i] 6 cyklů na warp • tyto detaily již nejsou nVidií publikovány (zjištěno měřením) • může se výrazně měnit s dalšími generacemi GPU, užitečné pro opravdu výkonově kritický kód Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti ooooooooooooooooooooo Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí ooooooo»o Pozor na sdílenou paměť U novějších GPU relativní rychlost sdílené paměti omezená • Fermi, Maxwell mají nižší propustnost ve srovnání s registry i v případě pouze jednoho operandu • Kepler je ještě více omezen, pokud nepřistupujeme k 64-bit číslům, používáme jen 1/2 přenosové kapacity Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOOOOO Překlad C for CUDA Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOO* Device kódy lze přeložit do PTX assembleru a binárních souborů • PTX je mezikód, neodpovídá přímo instrukcím prováděným na GPU • snáze se čte • hůře se zjišťuje, co se přesně na GPU děje Binární soubory lze deassemblovat pomocí nástroje cuobjdump • pro GT200 a novější • pro starší procesory decuda (produkt třetí strany, nemusí fungovat zcela správně) n > < & * 4 = > < = > s -00.0 Jiří Filipovič Výkon GPU hardware