Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí OOOOOOOOOOOOOOOOOOOO OOOOOOOOOOOOOOOOO OOOOOOOO Výkon GPU hardware Jiří Filipovič podzim 2013 Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti ♦OOOOOOOOOOOOOOOOOOO Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOO 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) • 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) Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí o«oooooooooooooooooo OOOOOOOOOOOOOOOOO OOOOOOOO 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 J128B aligned segment ................. Half warp of threads Optimalizace přístupu do paměti 00*00000000000000000 Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOO 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 Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti ooo»oooooooooooooooo Transpozice matic OOOOOOOOOOOOOOOOO Spojitý přístup do paměti (c.c. < 2.0) Rychlost instrukcí OOOOOOOO 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 0000*000000000000000 Transpozice matic OOOOOOOOOOOOOOOOO Spojitý přístup do paměti (c.c. < 2.0) Rychlost instrukcí OOOOOOOO Thready jsou zarovnané, blok elementů souvislý, pořadí není permutované - spojitý přístup na všech GPU. á i J 1 1 I 1 1 1 1 r t 1 1 1 1 Výkon GPU hardware Optimalizace přístupu do paměti OOOOO0OOOOOOOOOOOOOO Transpozice matic OOOOOOOOOOOOOOOOO Nezarovnaný přístup do paměti (c.c. < 2.0) Rychlost instrukcí OOOOOOOO 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ÄOOOOOOOOOOOOO Transpozice matic OOOOOOOOOOOOOOOOO Nezarovnaný přístup do paměti (c.c. < 2.0) Rychlost instrukcí OOOOOOOO Obdobný případ může vézt k použití dvou transakcí. —J ľ 71 m ttff 'li 77 ľi Vi Vi ] TÁ TEÉĽĽn Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti ooooooo»oooooooooooo Transpozice matic ooooooooooooooooo Rychlost instrukcí oooooooo Výkon při nezarovnaném přístupu (c.c. < 2.0) Starší GPU provádí pro každý element nejmenší možný přenos, tedy 32-bytů, což redukuje výkon na 1/8. Nové GPU (c.c. > 1.2) provádí dva přenosy. -W-GTX280 ? -*FX5600 2 4 6 8 10 12 14 16 Offset Optimalizace přístupu do paměti oooooooo«ooooooooooo Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOO 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 Transpozice matic Rychlost instrukcí OOOOOOOOO^OOOOOOOOOO OOOOOOOOOOOOOOOOO OOOOOOOO 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ÄOOOOOOOOO Transpozice matic OOOOOOOOOOOOOOOOO Přístup do globální paměti u Kepler (c.c. > 3.0) Rychlost instrukcí OOOOOOOO 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 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() Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti OOOOOOOOOOO^OOOOOOOO Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOO 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 oooooooooooo«ooooooo Transpozice matic ooooooooooooooooo HW organizace sdílené paměti Rychlost instrukcí OOOOOOOO 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 a 3.x 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 ooooooooooooo«oooooo Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOO 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, či přistupují k rozdílným podslovům 64-bitové banky u 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 Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti OOOOOOOOOOOOOOÍOOOOO Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOO Přístup bez konfliktů Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti ooooooooooooooo«oooo Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOO Vícecestné konflikty Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOÄOOO Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOO Broadcast Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti 00000000000000000*00 Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOO 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 3.x 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.x a 3.x 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]; Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOO0O Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOO 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 Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOOO* Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOO 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 oooooooooooooooooooo Transpozice matic ♦OOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOO 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 oooooooooooooooooooo Transpozice matic OÄOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOO Výkon Spustíme-li kód na GeForce GTX 280 s použitím dostatečně velké matice 4000 x 4000, bude propustnost 5.3 GB/s. Kde je problém? Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti oooooooooooooooooooo Transpozice matic OÄOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOO Výkon Spustíme-li kód na GeForce GTX 280 s použitím dostatečně velké matice 4000 x 4000, bude propustnost 5.3 GB/s. Kde je problém? Přístup do odata je prokládaný! Modifikujeme transpozici na kopírování: odata[y*n + x] = idata[y*n + x]; a získáme propustnost 112.4 GB/s. Pokud bychom přistupovali s prokládáním i k idata, bude výsledná rychlost 2.7 GB/s. Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOOOO Transpozice matic OO^OOOOOOOOOOOOOO Rychlost instrukcí OOOOOOOO 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í Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí oooooooooooooooooooo oo»oooooooooooooo oooooooo 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ě Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti Transpozice matic oooooooooooooooooooo ooo«ooooooooooooo Bloková transpozice Rychlost instrukcí OOOOOOOO __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) t ile[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]; 1 ► 1 -oc^o Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti oooooooooooooooooooo Transpozice matic OOOO^OOOOOOOOOOOO Rychlost instrukcí OOOOOOOO Výkon 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. Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti oooooooooooooooooooo Transpozice matic oooo»oooooooooooo Rychlost instrukcí OOOOOOOO Výkon 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 OOOOOOOOOOOOOOOOOOOO Transpozice matic OOOO^OOOOOOOOOOOO Rychlost instrukcí OOOOOOOO 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 Optimalizace přístupu do paměti oooooooooooooooooooo Transpozice matic oooo»oooooooooooo Rychlost instrukcí OOOOOOOO Výkon 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 • pokud v rámci bloků pouze kopírujeme, dosáhneme výkonu 94.9GB/s • něco ještě není optimální Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOOOO Transpozice matic OOOOOÄOOOOOOOOOOO Rychlost instrukcí OOOOOOOO 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]; Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOOOO Transpozice matic OOOOOÄOOOOOOOOOOO Rychlost instrukcí OOOOOOOO 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. Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOOOO Transpozice matic OOOOOÄOOOOOOOOOOO Rychlost instrukcí OOOOOOOO 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]; Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti oooooooooooooooooooo Transpozice matic 000000*0000000000 Rychlost instrukcí OOOOOOOO 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 (viz. partition camping) Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOOOO Transpozice matic OOOOOOOÄOOOOOOOOO Rychlost instrukcí OOOOOOOO Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí oooooooooooooooooooo oooooooo«oooooooo oooooooo 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í OOOOOOOOOOOOOOOOOOOO OOOOOOOOOÄOOOOOOO OOOOOOOO 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 oooooooooooooooooooo Transpozice matic OOOOOOOOOÄOOOOOOO Rychlost instrukcí OOOOOOOO Poklesy výkonu 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 Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti oooooooooooooooooooo Transpozice matic 0000000000*000000 Rychlost instrukcí OOOOOOOO 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! Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti oooooooooooooooooooo Transpozice matic ooooooooooo»ooooo Jak odstraníme partition camping Rychlost instrukcí OOOOOOOO 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 i -00.0 Výkon GPU hardware Optimalizace přístupu do paměti oooooooooooooooooooo Transpozice matic ooooooooooo»ooooo Jak odstraníme partition camping Rychlost instrukcí OOOOOOOO 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; Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti oooooooooooooooooooo Transpozice matic OOOOOOOOOOOOÄOOOO Rychlost instrukcí OOOOOOOO Výkon 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). Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOOOO Transpozice matic OOOOOOOOOOOOOÄOOO Rychlost instrukcí OOOOOOOO Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOOOO Transpozice matic OOOOOOOOOOOOOOÄOO Rychlost instrukcí OOOOOOOO Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti oooooooooooooooooooo Transpozice matic OOOOOOOOOOOOOOO^O Rychlost instrukcí OOOOOOOO 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... Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti oooooooooooooooooooo Transpozice matic OOOOOOOOOOOOOOOO* Rychlost instrukcí OOOOOOOO 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 oooooooooooooooooooo Provádění instrukcí Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí •OOOOOOO 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 OOOOOOOOOOOOOOOOOOOO Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí OÄOOOOOO 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 oooooooooooooooooooo Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí 00*00000 Aritmetické operace Operace s plovoucí řádovou čárkou (propustnost na MP) • sčítání, násobení 8 (1.x), 32 (2.0), 48 (2.1), 192 (3.x) • násobení a sčítání může být u c.c. 1.x kombinováno dojedná instrukce MAD o nižší přesnost • rychlost 1 cyklus na SP • —fadd-rn() 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 1/8 (1.3), 1/2 (2.0), 1/12 (2.1), 1/24 (3.0), 1/3 (3.5) • převrácená hodnota 2 (1.x), 4 (2.0), 8 (2.1), 32 (3.x) • dělení relativně pomalé (u c.c. 1.x v průměru cca 1.23) • rychlejší varianta pomocí __fdividef(x, y) 1.6 (c.c. 1.x) • reciproká druhá odmocnina 2 (1.x), 4 (2^.0)^(2.1^ 3g (3.^) Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí OOOOOOOOOOOOOOOOOOOO OOOOOOOOOOOOOOOOO 000*0000 Aritmetické operace Operace s plovoucí řádovou čárkou • _.sinf(x), -cosffx), -expf(x) 2 (c.c. 1.x), 4 (c.c. 2.0), 8 (c.c 1.2), 32 (3.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 (160 u c.c. 3.0) • násobení u c.c. 1.x 2 instrukce na MP • __mu/24(x, y) a __umul24(x, y) 8 instrukcí • násobení u c.c. 2.x 16, u c.c. 3.x 32 instrukcí na MP, 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) m -OQ.O Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti OOOOOOOOOOOOOOOOOOOO Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí 0000*000 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 OOOOOOOOOOOOOOOOOOOO Ostatní operace Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí OOOOOÄOO 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 oooooooooooooooooooo Transpozice matic ooooooooooooooooo Rychlost instrukcí 000000*0 Pozor na sdílenou paměť Pokud nedojde ke konfliktům bank, je sdílená paměť rychlá téměř 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 OOOOOOOOOOOOOOOOOOOO Překlad C for CUDA Transpozice matic OOOOOOOOOOOOOOOOO Rychlost instrukcí 0000000» 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