Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí ooooooooooooooooooo ooooooooooooooooo oooooooo Výkon GPU hardware Jiří Filipovič podzim 2010 Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti •oooooooooooooooooo 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) • 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) □ g - = = -^c^O Jiří Filipouič Výkon GPU hardware Optimalizace přístupu do paměti o»ooooooooooooooooo Transpozice matic ooooooooooooooooo Rychlost instrukcí 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 > 128B aligned segment ......I.......... Half warp of threads Výkon GPU hardware Optimalizace přístupu do paměti oo»oooooooooooooooo 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 Optimalizace přístupu do paměti ooo»ooooooooooooooo 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 nejnižší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 oooo»oooooooooooooo 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. ■ k A 1 1 ! 1 1 1 1 1 I 1 1 1 1 1 Výkon GPU hardware Optimalizace přístupu do paměti ooooo»ooooooooooooo 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. iVi'i'i'i'i'i'i'i'i'iŕŤWn Výkon GPU hardware Optimalizace přístupu do paměti oooooo»oooooooooooo Transpozice matic ooooooooooooooooo Nezarovnaný prístup do paměti(c.c. < 2.0) Rychlost instrukcí oooooooo Obdobný případ může vézt k nutnosti použít dvě transakce. □ □ - - Jiří Filipovie Výkon GPU hardware Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí ooooooo»ooooooooooo ooooooooooooooooo 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. 140 oH— D 2 4 B B 10 12 14 16 Offset Optimalizace přístupu do paměti oooooooo»oooooooooo 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-1—li i—i 0 2 4 B 8 10 12 14 16 18 Stride □ - = = ^Q^O Jiří Filipouič Výkon GPU hardware Optimalizace přístupu do paměti ooooooooo«ooooooooo Transpozice matic ooooooooooooooooo Rychlost instrukcí oooooooo 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) Jiří Filipouič Výkon GPU hardware Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí oooooooooo»oooooooo ooooooooooooooooo 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, špatně viditelné při jemnozrnném pohledu □ - = = ^Q^O Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti ooooooooooo«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 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 Optimalizace přístupu do paměti oooooooooooo»oooooo Konflikty bank Transpozice matic ooooooooooooooooo Rychlost instrukcí oooooooo 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ézt 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 Optimalizace přístupu do paměti ooooooooooooo«ooooo Prístup bez konfliktů Transpozice matic ooooooooooooooooo Rychlost instrukcí oooooooo Optimalizace přístupu do paměti oooooooooooooo»oooo Transpozice matic ooooooooooooooooo Rychlost instrukcí oooooooo Vřcecestné konflikty Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti ooooooooooooooo«ooo Transpozice matic ooooooooooooooooo Rychlost instrukcí oooooooo Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti oooooooooooooooo»oo Vzory přístupu Transpozice matic ooooooooooooooooo Rychlost instrukcí oooooooo Zarovnání není třeba, negeneruje bank conflicty int x = s[threadldx.x + offset]; Prokládání negeneruje konflikty, je-li c liché int x = s[threadldx.x * c]; Přístup ke stejné proměnné negeneruje na c.c. 2.x konflikty nikdy, na l,x je-li počet c threadů přistupující k proměnné násobek 16 int x = s[threadldx.x % c]; Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí oooooooooooooooooso ooooooooooooooooo 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 0OOOOOOOOOOOOOOOOO» 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 □ - = = ^Q^O Jiří Filipouič Výkon GPU hardware Optimalizace přístupu do paměti ooooooooooooooooooo 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]; Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí ooooooooooooooooooo o»ooooooooooooooo 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? Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí ooooooooooooooooooo o»ooooooooooooooo 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. Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí ooooooooooooooooooo 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í Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí ooooooooooooooooooo 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ě □ - = = ^Q^O Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti ooooooooooooooooooo Transpozice matic ooo»ooooooooooooo Rychlost instrukcí oooooooo 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 ooooooooooooooooooo Transpozice matic oooo»oooooooooooo Rychlost instrukcí oooooooo 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 ooooooooooooooooooo Transpozice matic oooo»oooooooooooo Rychlost instrukcí oooooooo 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 ooooooooooooooooooo Transpozice matic oooo»oooooooooooo Rychlost instrukcí oooooooo 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 ooooooooooooooooooo Transpozice matic oooo»oooooooooooo Rychlost instrukcí oooooooo 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 Optimalizace přístupu do paměti ooooooooooooooooooo Sdílená paměť Transpozice matic ooooo»ooooooooooo Rychlost instrukcí oooooooo 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 Optimalizace přístupu do paměti ooooooooooooooooooo Sdílená paměť Transpozice matic ooooo»ooooooooooo Rychlost instrukcí oooooooo 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 Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí ooooooooooooooooooo ooooo»ooooooooooo 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]; Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí ooooooooooooooooooo oooooosoooooooooo 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) Optimalizace přístupu do paměti ooooooooooooooooooo Transpozice matic ooooooo»ooooooooo Rychlost instrukcí oooooooo Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí ooooooooooooooooooo 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 ooooooooooooooooooo Poklesy výkonu Transpozice matic ooooooooo»ooooooo Rychlost instrukcí oooooooo Pro některé velikosti problému výkon klesá, v tomto chování lze nalézt regularitu s Výkon GPU hardware Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí ooooooooooooooooooo ooooooooo»ooooooo oooooooo Poklesy výkonu Pro některé velikosti problému výkon klesá, v tomto chování lze nalézt regularitu • u matic o velikosti dělitelné 512 dosahujeme pouze cca 19 GB/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 ooooooooooooooooooo Transpozice matic oooooooooo»oooooo 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! Optimalizace přístupu do paměti ooooooooooooooooooo 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 s Výkon GPU hardware Optimalizace přístupu do paměti ooooooooooooooooooo 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 blockldx.x = (blockldx.x+blockldx.y) % gridDim.x; □ - = = ^Q^O Jiří Filipouič Výkon GPU hardware Optimalizace přístupu do paměti ooooooooooooooooooo Transpozice matic oooooooooooo»oooo Rychlost instrukcí oooooooo 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). s Výkon GPU hardware Optimalizace přístupu do paměti ooooooooooooooooooo Transpozice matic ooooooooooooo»ooo Rychlost instrukcí oooooooo Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti ooooooooooooooooooo Transpozice matic oooooooooooooo»oo Rychlost instrukcí oooooooo Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti ooooooooooooooooooo Zhodnocení výkonu Transpozice matic ooooooooooooooo«o Rychlost instrukcí oooooooo 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ý sekvenční algoritmus... Optimalizace přístupu do paměti ooooooooooooooooooo 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í □ g - = = -0*3.0 Transpozice matic 0000000000000000» Rychlost instrukci oooooooo Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti ooooooooooooooooooo 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 optimální kód Jiří Filipouič Výkon GPU hardware Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí ooooooooooooooooooo ooooooooooooooooo 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 Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí ooooooooooooooooooo ooooooooooooooooo oosooooo Aritmetické operace Operace s plovoucí řádovou čárkou (propustnost na MP) • sčítání, násobení 8 (1.x), 32 (2.0), 48 (2.1) • násobení a sčítání může být u c.c. 1.x kombinováno dojedná instrukce MAD • 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) • převrácená hodnota 2 (1.x), 4 (2.0) a 8 (2.1) • 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) a 8 (2.1) » konverze typů 8 (c.c. 1.x), 16 (c.c. 2.x) Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti Transpozice matic Rychlost instrukcí ooooooooooooooooooo ooooooooooooooooo ooo»oooo Aritmetické operace Operace s plovoucí řádovou čárkou • __sinf(x), -cosf(x), __expf(x) 2 (c.c. 1.x), 4 (c.c. 2.0), 8 (c.c. 1.2) » 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 2 instrukce na M P • —mul24(x, y) a __umul24(x, y) 8 instrukcí • násobení u c.c. 2.x stejně rychlé jako u plovoucí řádové čárky, 24-bitová verze naopak pomalá o dělení a modulo velmi pomalé, pokud je n mocnina 2, můžeme využít • i/n je ekvivalentní ; >> íog2{n) • i%n je ekvivalentní i&í(n — 1) Jiří Filipovič Výkon GPU hardware Optimalizace přístupu do paměti ooooooooooooooooooo Smyčky Transpozice matic ooooooooooooooooo Rychlost instrukcí oooo»ooo 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 Optimalizace přístupu do paměti ooooooooooooooooooo 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í • bitové operace • 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 :-)) □ - = = -0*3*0 Jiří Filipouič Výkon GPU hardware Optimalizace přístupu do paměti ooooooooooooooooooo Transpozice matic ooooooooooooooooo Rychlost instrukcí oooooo»o Pozor na sdílenou paměť Pokud nedojde ke konfliktům bank, je sdílená paměť 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 ooooooooooooooooooo Transpozice matic ooooooooooooooooo Rychlost instrukcí 0000000» Překlad C for CUDA 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 a překladač do nativního GPU kódu má být uvolněn Binární soubory lze deassemblovat pomocí nástroje decuda • produkt třetí strany • nemusí fungovat zcela správně • přesto dosti užitečný Jiří Filipouič Výkon GPU hardware