Optimalizace míst Jiří Filipovič podzim 2011 Jiří Filipovič Optimalizace Redukce Obecné zásady OOOOOOOOOOOOOOOOOOO OOOOOOO Revize násobení matic •oooooooooooooo Naivní implementace Hledání slabých míst OOOOOOOO __global__ void mmul(float *A, float *B, float *C, int n){ int x = blockldx.x*blockDim.x + threadldx.x; int y = blockldx.y*blockDim.y + threadldx.y; float tmp = 0; for (int k = 0; k < n; k++) tmp += A[y*n+k] * B[k*n+x]; C[y*n + x] — tmp; } 4fiP> -š -00,0 Jiří Filipovič Optimalizace Revize násobení matic Redukce Obecné zásady Hledání slabých míst OÄOOOOOOOOOOOOO ooooooooooooooooooo ooooooo oooooooo Co jsme se naučili Naivní implementace algoritmu • každý thread zpracovává odděleně jeden element výsledné matice • omezena propustností paměti • teoretické maximum jsme určili jako 66.8GFIops • výkon velmi závislý na uspořádání threadů - bloky 128 x 1 dávají výkon 36.6GFIops, bloky 1 x 128 3.9GFIops Revize násobení matic OÄOOOOOOOOOOOOO Co jsme se naučili Naivní implementace algoritmu • každý thread zpracovává odděleně jeden element výsledné matice • omezena propustností paměti • teoretické maximum jsme určili jako 66.8GFIops • výkon velmi závislý na uspořádání threadů - bloky 128 x 1 dávají výkon 36.6GFIops, bloky 1 x 128 3.9GFIops Nyní rozumíme rozdílným výsledkům • teoretického maxima nelze docílit - z paměti GPU přenášíme po nejméně 32-bytových částech, musíme tedy přenést více dat, než je nutné • je-li 128 threadů v bloku zarovnáno ve smyslu osy x, je přenos dat neprokládaný, v opačném případě je prokládaný Redukce Obecné zásady OOOOOOOOOOOOOOOOOOO OOOOOOO Hledání slabých míst OOOOOOOO Jiří Filipovič Optimalizace Revize násobení matic Redukce Obecné zásady Hledání slabých míst 00*000000000000 ooooooooooooooooooo ooooooo oooooooo Co jsme se naučili Navrhli jsme blokovou implementaci • každý blok threadů načítá bloky matic A a B do sdílené paměti, znovuužívá data ke snížení omezení přenosovou rychlostí globální paměti • teoretické maximum 568GFIops, dosáhli jsme 198GFIops S novými znalostmi můžeme jeho implementaci přehodnotit... Revize násobení matic Redukce Obecné zásady Hledání slabých míst ooo»ooooooooooo ooooooooooooooooooo ooooooo oooooooo Násobení po blocích _global__ void mmul(float *A, float *B, float *C, int n){ int bx = blockldx.x; int by = blockldx.y; int tx = threadldx.x; int ty = threadldx.y; __shared__ float As[BLOCK][BLOCK]; __shared__ float Bs[BLOCK][BLOCK]; float Csub = O.Of; for (int b = 0; b < n/BLOCK; b++){ As[ty][tx] — A[(ty + by*BL0CK)*n + b*BL0CK+tx]; Bs[ty][tx] — B[(ty + b*BL0CK)*n + bx*BL0CK+tx]; __syncthreads(); for (int k = 0; k < BLOCK; k++) Csub += As[ty][k]*Bs[k][tx]; __syncthreads(); } C[(ty + by*BL0CK)*n + bx*BL0CK+tx] = Csub; } Jiří Filipovič Optimalizace Revize násobení matic oooo«oooooooooo Redukce Obecné zásady ooooooooooooooooooo ooooooo Hledání slabých míst oooooooo Hříchy implementace As[ty][tx] — A[(ty + by*BL0CK)*n + b*BL0CK+tx]; Bsjtyjjtx] — B[(ty + b*BL0CK)*n + bx*BLOCK+tx]; C[(ty + by*BL0CK)*n + bx*BLOCK+tx] = Csub; Přístup do globální paměti se zdá být v pořádku. Csub += As[ty][k]*Bs[k][tx]; Přístup do sdílené také • má-li blok threadů velikost ve smyslu osy x násobek velikosti warpu, dochází u proměnné As k broadcastu • proměnná Bs je čtena v souvislých řádcích, přístup tedy negeneruje konflikty bank Revize násobení matic Redukce Obecné zásady Hledání slabých míst ooooo»ooooooooo ooooooooooooooooooo ooooooo oooooooo Teoretické maximum Lze určit přesněji teoretické omezení výkonu? • maximum jsme určili podle výkonu GPU v MAD instrukcích (622GFIops) • nyní víme, že MAD instrukce pracující s operandem ve sdílené paměti pracují rychlostí 6 taktů na warp • nově lze tedy teoretické maximum určit jako 415GFIops • stále jsme od něj však daleko Revize násobení matic Redukce Obecné zásady Hledání slabých míst oooooo«oooooooo ooooooooooooooooooo ooooooo oooooooo Ztráty výkonu Co nás vzdaluje od maxima? • overhead spuštění kernelu a spouštění threadů • z principu se mu nevyhneme, počet threadů lze redukovat • operace „režije" • pointerová aritmetika, cykly • lze redukovat • synchronizace • může a nemusí být problém • load/store ve výpočtu • dva operandy v SMEM na jednu MAD instrukci • je tedy zapotřebí jeden load na jednu MAD Počítáme-li výkonový strop pro kombinaci load + MAD s operandem ve sdílené paměti, dostaneme se k omezení 244 GFIops. • od toho již nejsou naměřené výsledky příliš vzdáleny Jiří Filipovič Optimalizace Revize násobení matic ooooooo»ooooooo Redukce Obecné zásady ooooooooooooooooooo ooooooo Hledání slabých míst oooooooo Nalezení lepší implementace Lze počet load instrukcí omezit? Hledání slabých míst oooooooo Nalezení lepší implementace Lze počet load instrukcí omezit? • data ve sdílené paměti snižují přenosy z paměti globální i -00.0 Hledání slabých míst oooooooo Nalezení lepší implementace Lze počet load instrukcí omezit? • data ve sdílené paměti snižují přenosy z paměti globální • můžeme snížit přenosy ze sdílené paměti pomocí dat v registrech? i -00.0 Revize násobení matic ooooooo»ooooooo Redukce Obecné zásady ooooooooooooooooooo ooooooo Hledání slabých míst oooooooo Nalezení lepší implementace Lze počet load instrukcí omezit? • data ve sdílené paměti snižují přenosy z paměti globální • můžeme snížit přenosy ze sdílené paměti pomocí dat v registrech? • můžeme - stačí nechat pracovat méně threadů nad více daty Jiří Filipovič Optimalizace Revize násobení matic ooooooo»ooooooo Redukce Obecné zásady ooooooooooooooooooo ooooooo Hledání slabých míst oooooooo Nalezení lepší implementace Lze počet load instrukcí omezit? • data ve sdílené paměti snižují přenosy z paměti globální • můžeme snížit přenosy ze sdílené paměti pomocí dat v registrech? • můžeme - stačí nechat pracovat méně threadů nad více daty Blok o velikosti m x n threadů necháme pracovat s daty o velikosti m x m, kde m = n ■ k; k G N. • větší bloky potenciálně nevýhodné kvůli synchronizaci • menší bloky potenciálně nevýhodné kvůli nižšímu dosažitelnému paralelismu a overheadu daném pointerovou aritmetikou • experimentálně najdeme vhodnou velikost bloku Jiří Filipovič Optimalizace Revize násobení matic Redukce Obecné zásady Hledání slabých míst oooooooo«oooooo ooooooooooooooooooo ooooooo oooooooo Nalezení lepší implementace Nejlepší výsledky dosaženy pro bloky matice 32 x 32, na kterých pracují bloky 32 x 16 threadů. • půl loadu na jednu MAD instrukci dává teoretické omezení 311GFIops • naměřili jsme 235.4 GFIops • něco je ještě špatně Revize násobení matic ooooooooo»ooooo Redukce Obecné zásady ooooooooooooooooooo ooooooo Hledání slabých míst oooooooo Deassembling kódu Zaměříme se na vnitřní smyčku Csubl += As[ty][k]*Bs[k] [tx] ; Csub2 += As [ ty + 16] [k]* Bs [k ] [ tx ] ; mov.b32 $r0, s[$ofs4+0x0000] add.b32 $ofs4, $ofs2, 0x00000180 mad.rn.f32 $r7, s[$ofs 1+0x0008] , $r0 , $r7 mad.rn.f32 $r8, s[$ofs3+0x0008j, $r0, $r8 i -00.0 Revize násobení matic ooooooooo»ooooo Redukce Obecné zásady ooooooooooooooooooo ooooooo Hledání slabých míst oooooooo Deassembling kódu Zaměříme se na vnitřní smyčku Csubl += As[ty][k]*Bs[k] [tx] ; Csub2 += As [ ty + 16] [k]* Bs [k ] [ tx ] ; mov.b32 $r0, s[$ofs4+0x0000] add.b32 $ofs4, $ofs2, 0x00000180 mad.rn.f32 $r7, s[$ofs 1+0x0008] , $r0 , $r7 mad.rn.f32 $r8, s[$ofs3+0x0008j, $r0, $r8 Kompilátor dokázal převést adresaci přes k na konstantní offsety pouze u proměnné As 9 k Bs je přistupováno prokládané • znamená to jednu add instrukci navíc Jiří Filipovič Optimalizace Revize násobení matic Redukce Obecné zásady Hledání slabých míst oooooooooo«oooo ooooooooooooooooooo ooooooo oooooooo Odstranění add instrukce Do pole Bs můžeme ukládat transponovaná data, pak vypadá kód vnitřní smyčky takto Csubl += As[ty][k]*Bs[tx] [k] ; Csub2 += As [ty + 16][k]*Bs [tx] [k] ; Revize násobení matic oooooooooo«oooo Redukce Obecné zásady ooooooooooooooooooo ooooooo Hledání slabých míst oooooooo Odstranění add instrukce Do pole Bs můžeme ukládat transponovaná data, pak vypadá kód vnitřní smyčky takto Csubl += As[ty][k]*Bs[tx] [k] ; Csub2 += As [ty + 16][k]*Bs [tx] [k] ; Ve výsledném assembleru již instrukce add chybí mov.b32 $r0, s[$ofs4+0x0008] mad.rn.f32 $r6, s[$ofs3+0x0034], $r0, $r6 mad.rn.f32 $r8, s[$ofs 1+0x0008 j , $r0 , $r8 Revize násobení matic oooooooooo«oooo Redukce Obecné zásady ooooooooooooooooooo ooooooo Hledání slabých míst oooooooo Odstranění add instrukce Do pole Bs můžeme ukládat transponovaná data, pak vypadá kód vnitřní smyčky takto Csubl += As[ty][k]*Bs[tx] [k] ; Csub2 += As [ty + 16][k]*Bs [tx] [k] ; Ve výsledném assembleru již instrukce add chybí mov.b32 $r0, s[$ofs4+0x0008] mad.rn.f32 $r6, s[$ofs3+0x0034], $r0, $r6 mad.rn.f32 $r8, s[$ofs 1+0x0008 j , $r0 , $r8 Nový problém - konflikty bank sdílené paměti Revize násobení matic oooooooooo«oooo Redukce Obecné zásady ooooooooooooooooooo ooooooo Hledání slabých míst oooooooo Odstranění add instrukce Do pole Bs můžeme ukládat transponovaná data, pak vypadá kód vnitřní smyčky takto Csubl += As[ty][k]*Bs[tx] [k] ; Csub2 += As [ty + 16][k]*Bs [tx] [k] ; Ve výsledném assembleru již instrukce add chybí mov.b32 $r0, s[$ofs4+0x0008] mad.rn.f32 $r6, s[$ofs3+0x0034], $r0, $r6 mad.rn.f32 $r8, s[$ofs 1+0x0008 j , $r0 , $r8 Nový problém - konflikty bank sdílené paměti • vyřeší padding Revize násobení matic oooooooooo«oooo Redukce Obecné zásady ooooooooooooooooooo ooooooo Hledání slabých míst oooooooo Odstranění add instrukce Do pole Bs můžeme ukládat transponovaná data, pak vypadá kód vnitřní smyčky takto Csubl += As[ty][k]*Bs[tx] [k] ; Csub2 += As [ty + 16][k]*Bs [tx] [k] ; Ve výsledném assembleru již instrukce add chybí mov.b32 $r0, s[$ofs4+0x0008] mad.rn.f32 $r6, s[$ofs3+0x0034], $r0, $r6 mad.rn.f32 $r8, s[$ofs 1+0x0008 j , $r0 , $r8 Nový problém - konflikty bank sdílené paměti • vyřeší padding Výsledná rychlost: 276.2 GFIops. Jiří Filipovič Optimalizace Redukce Obecné zásady OOOOOOOOOOOOOOOOOOO OOOOOOO Revize násobení matic OOOOOOOOOOO^OOO Lze matice násobit ještě rychleji? Hledání slabých míst OOOOOOOO Naměřený výkon je již poměrně blízký teoretickému maximu • rozdíl je dán spouštěním kernelu/threadů, synchronizací a pointerovou aritmetikou • chceme-li dosáhnout vyšší rychlosti, je třeba přehodnotit algoritmus Jiří Filipovič Optimalizace Redukce Obecné zásady OOOOOOOOOOOOOOOOOOO OOOOOOO Revize násobení matic OOOOOOOOOOO^OOO Lze matice násobit ještě rychleji? Hledání slabých míst OOOOOOOO Naměřený výkon je již poměrně blízký teoretickému maximu • rozdíl je dán spouštěním kernelu/threadů, synchronizací a pointerovou aritmetikou • chceme-li dosáhnout vyšší rychlosti, je třeba přehodnotit algoritmus Zásadním problémem je, že spolu násobíme dvě matice ve sdílené paměti • nutnost provádět load instrukce spolu s MAD instrukcemi Jiří Filipovič Optimalizace Redukce Obecné zásady OOOOOOOOOOOOOOOOOOO OOOOOOO Revize násobení matic OOOOOOOOOOO^OOO Lze matice násobit ještě rychleji? Hledání slabých míst OOOOOOOO Naměřený výkon je již poměrně blízký teoretickému maximu • rozdíl je dán spouštěním kernelu/threadů, synchronizací a pointerovou aritmetikou • chceme-li dosáhnout vyšší rychlosti, je třeba přehodnotit algoritmus Zásadním problémem je, že spolu násobíme dvě matice ve sdílené paměti • nutnost provádět load instrukce spolu s MAD instrukcemi Můžeme mít ve sdílené paměti jen jeden blok? Jiří Filipovič Optimalizace Revize násobení matic OOOOOOOOOOOOÄOO Redukce Obecné zásady OOOOOOOOOOOOOOOOOOO OOOOOOO Hledání slabých míst OOOOOOOO Přehodnocený blokový přístup Namísto čtvercových bloků v matici C můžeme použít obdélníkové Jiří Filipovič Optimalizace Hledání slabých míst OOOOOOOO Přehodnocený blokový přístup Namísto čtvercových bloků v matici C můžeme použít obdélníkové • provádíme iterativně rank-1 update bloků v C ze sloupce matice A a řádku matice B i -00.0 Hledání slabých míst OOOOOOOO Přehodnocený blokový přístup Namísto čtvercových bloků v matici C můžeme použít obdélníkové • provádíme iterativně rank-1 update bloků v C ze sloupce matice A a řádku matice B • sloupce je nutno číst se sdílené paměti i -00.0 Revize násobení matic Redukce Obecné zásady Hledání slabých míst OOOOOOOOOOOOÄOO ooooooooooooooooooo ooooooo oooooooo Přehodnocený blokový přístup Namísto čtvercových bloků v matici C můžeme použít obdélníkové • provádíme iterativně rank-1 update bloků v C ze sloupce matice A a řádku matice B • sloupce je nutno číst se sdílené paměti • řádky můžeme načítat postupně, lze tedy použít data v registrech Revize násobení matic Redukce Obecné zásady Hledání slabých míst OOOOOOOOOOOOÄOO ooooooooooooooooooo ooooooo oooooooo Přehodnocený blokový přístup Namísto čtvercových bloků v matici C můžeme použít obdélníkové • provádíme iterativně rank-1 update bloků v C ze sloupce matice A a řádku matice B • sloupce je nutno číst se sdílené paměti • řádky můžeme načítat postupně, lze tedy použít data v registrech • výsledný blok může být uložen v registrech Revize násobení matic Redukce Obecné zásady Hledání slabých míst OOOOOOOOOOOOÄOO ooooooooooooooooooo ooooooo oooooooo Přehodnocený blokový přístup Namísto čtvercových bloků v matici C můžeme použít obdélníkové • provádíme iterativně rank-1 update bloků v C ze sloupce matice A a řádku matice B • sloupce je nutno číst se sdílené paměti • řádky můžeme načítat postupně, lze tedy použít data v registrech • výsledný blok může být uložen v registrech • pracujeme tedy pouze s jedním operandem ve sdílené paměti, není nutný load Revize násobení matic Redukce Obecné zásady Hledání slabých míst OOOOOOOOOOOOÄOO ooooooooooooooooooo ooooooo oooooooo Přehodnocený blokový přístup Namísto čtvercových bloků v matici C můžeme použít obdélníkové • provádíme iterativně rank-1 update bloků v C ze sloupce matice A a řádku matice B • sloupce je nutno číst se sdílené paměti • řádky můžeme načítat postupně, lze tedy použít data v registrech • výsledný blok může být uložen v registrech • pracujeme tedy pouze s jedním operandem ve sdílené paměti, není nutný load • není nutná aritmetika uprostřed smyčky (viz předchozí optimalizace) Revize násobení matic Redukce Obecné zásady Hledání slabých míst OOOOOOOOOOOOÄOO ooooooooooooooooooo ooooooo oooooooo Přehodnocený blokový přístup Namísto čtvercových bloků v matici C můžeme použít obdélníkové • provádíme iterativně rank-1 update bloků v C ze sloupce matice A a řádku matice B • sloupce je nutno číst se sdílené paměti • řádky můžeme načítat postupně, lze tedy použít data v registrech • výsledný blok může být uložen v registrech • pracujeme tedy pouze s jedním operandem ve sdílené paměti, není nutný load • není nutná aritmetika uprostřed smyčky (viz předchozí optimalizace) • teoretické maximum výkonu je tak omezeno rychlostí instrukce MAD pracující se sdílenou pamětí na cca 415GFIops Jiří Filipovič Optimalizace Revize násobení matic ooooooooooooo«o Implementace Redukce Obecné zásady OOOOOOOOOOOOOOOOOOO OOOOOOO Hledání slabých míst OOOOOOOO Nejvyšší rychlosti bylo dosaženo s konfigurací • matice A zpracovávána po blocích 16 x 16, uložených ve sdílené paměti • matice B zpracovávána po blocích 64 x 1, uložených v registrech • bloky matice C mají tedy rozměr 64 x 16, jsou uloženy v registrech Jiří Filipovič Optimalizace Revize násobení matic ooooooooooooo«o Implementace Redukce Obecné zásady OOOOOOOOOOOOOOOOOOO OOOOOOO Hledání slabých míst OOOOOOOO Nejvyšší rychlosti bylo dosaženo s konfigurací • matice A zpracovávána po blocích 16 x 16, uložených ve sdílené paměti • matice B zpracovávána po blocích 64 x 1, uložených v registrech • bloky matice C mají tedy rozměr 64 x 16, jsou uloženy v registrech Dosažená rychlost této implementace 375GFIops. Jiří Filipovič Optimalizace Revize násobení matic 00000000000000« Redukce Obecné zásady ooooooooooooooooooo ooooooo Hledání slabých míst OOOOOOOO Shrnutí Implementace rychlost rel. A abs. A Naivní implementace, thready 1 x 128 3.9 GFIops Naivní implementace 36.6 GFIops 9.4x 9.4x Blokový přístup 198GFIops 5.4x 51x Bloky 32 x 16 pracující s daty 32 x 16 235 GFIops 1.19x 60 x Odstranění ADD instrukce 276 GFIops 1.17x 71x Jen jeden blok ve sdílené paměti 375 GFIops 1.36x 96 x Jiří Filipovič Optimalizace Revize násobení matic 00000000000000« Redukce Obecné zásady ooooooooooooooooooo ooooooo Hledání slabých míst OOOOOOOO Shrnutí Implementace rychlost rel. A abs. A Naivní implementace, thready 1 x 128 3.9 GFIops Naivní implementace 36.6 GFIops 9.4x 9.4x Blokový přístup 198GFIops 5.4x 51x Bloky 32 x 16 pracující s daty 32 x 16 235 GFIops 1.19x 60 x Odstranění ADD instrukce 276 GFIops 1.17x 71x Jen jeden blok ve sdílené paměti 375 GFIops 1.36x 96 x • Nejzásadnější je redukce poměru aritmetických operací k paměťovým přenosům a základní optimalizace přístupu do paměti. Jiří Filipovič Optimalizace Revize násobení matic Redukce Obecné zásady Hledání slabých míst 00000000000000« ooooooooooooooooooo ooooooo oooooooo Shrnutí Implementace rychlost rel. A abs. A Naivní implementace, thready 1 x 128 3.9 GFIops Naivní implementace 36.6 GFIops 9.4x 9.4x Blokový přístup 198GFIops 5.4x 51x Bloky 32 x 16 pracující s daty 32 x 16 235 GFIops 1.19x 60 x Odstranění ADD instrukce 276 GFIops 1.17x 71x Jen jeden blok ve sdílené paměti 375 GFIops 1.36x 96 x • Nejzásadnější je redukce poměru aritmetických operací k paměťovým přenosům a základní optimalizace přístupu do paměti. • Optimalizace na úrovni instrukcí je relativně náročná, avšak pro kritické kódy může přinést relativně významné zrychlení. Jiří Filipovič Optimalizace Redukce Obecné zásady #000000000000000000 OOOOOOO Revize násobení matic OOOOOOOOOOOOOOO Součet prvků vektoru Hledání slabých míst OOOOOOOO Pro vektor v o n prvcích chceme spočítat x = Y^i=i v'f- Jiří Filipovič Optimalizace Redukce Obecné zásady #000000000000000000 OOOOOOO Revize násobení matic OOOOOOOOOOOOOOO Součet prvků vektoru Hledání slabých míst OOOOOOOO Pro vektor von prvcích chceme spočítat x = Yľi=i vi-Zápis (hloupý) v jazyce C int x = 0; for (int i — 0; i < n; i++) x += v[i ] ; Jednotlivé iterace cyklu jsou na sobě závislé. Revize násobení matic Redukce Obecné zásady Hledání slabých míst OOOOOOOOOOOOOOO »000000000000000000 ooooooo oooooooo Součet prvků vektoru Pro vektor v o n prvcích chceme spočítat x = Y11=i vi-Zápis (hloupý) v jazyce C int x = O; for (int i — 0; i < n; i++) x += v[i ] ; Jednotlivé iterace cyklu jsou na sobě závislé. • nemůžeme udělat všechnu práci paralelně • sčítání je však (alespoň teoreticky :-)) asocitativní • není tedy nutno počítat sekvenčně Revize násobení matic ooooooooooooooo Redukce Obecné zásady o«ooooooooooooooooo ooooooo Hledání slabých míst oooooooo Paralelní algoritmus Představený sekvenční algoritmus provádí pro 8 prvků výpočet: ((((((1/1 + v2) + v3) + 1/4) + 1/5) + v6) + vr) + v8 Jiří Filipovič Optimalizace Revize násobení matic Redukce Obecné zásady Hledání slabých míst ooooooooooooooo o«ooooooooooooooooo ooooooo oooooooo Paralelní algoritmus Představený sekvenční algoritmus provádí pro 8 prvků výpočet: ((((((1/1 + v2) + 1/3) + 1/4) + 1/5) + v6) + vr) + v8 Sčítání je asociativní... spřeházejme tedy závorky: {{vi + v2) + (1/3 + 1/4)) + ((1/5 + v6) + {vj + v8)) Revize násobení matic ooooooooooooooo Redukce Obecné zásady o«ooooooooooooooooo ooooooo Hledání slabých míst oooooooo Paralelní algoritmus Představený sekvenční algoritmus provádí pro 8 prvků výpočet: ((((((1/1 + v2) + v3) + 1/4) + 1/5) + v6) + vr) + v8 Sčítání je asociativní... spřeházejme tedy závorky: {{vi + v2) + (1/3 + 1/4)) + ((1/5 + v6) + {vj + v8)) Nyní můžeme pracovat paralelně • v prvním kroku provedeme 4 sčítání • ve druhém dvě • ve třetím jedno Celkově stejné množství práce (n — 1 sčítání), ale v log2 n paralelních krocích! Jiří Filipovič Optimalizace Revize násobení matic Redukce Obecné zásady Hledání slabých míst ooooooooooooooo oo»oooooooooooooooo ooooooo oooooooo Paralelní algoritmus Našli jsme vhodný paralelní algoritmus • provádí stejné množství operací jako sériová verze • při dostatku procesorů je proveden v logaritmickém čase Sčítáme výsledky předešlých součtů • předešlé součty provádělo více threadů • vyžaduje globální bariéru Jiří Filipovič Optimalizace Revize násobení matic Redukce Obecné zásady Hledání slabých míst OOOOOOOOOOOOOOO OOOÄOOOOOOOOOOOOOOO ooooooo oooooooo Naivní přístup Nejjednodušší schéma algoritmu: • kernel pro sudá i < n provede v[i] += v[i+l] • opakujeme pro n /= 2 dokud n > 1 Omezení výkonu • 2n čtení z globální paměti • n zápisů do globální paměti • log2 n volání kernelu Na jednu aritmetickou operaci připadají 3 paměťové přenosy, navíc je nepříjemný overhead spouštění kernelu. Revize násobení matic Redukce Obecné zásady Hledání slabých míst ooooooooooooooo oooo»oooooooooooooo ooooooo oooooooo Využití rychlejší paměti V rámci volání kernelu můžeme posčítat více, než jen dvojice • každý blok bx načte m prvků do sdílené paměti • provede redukci (ve sdílené paměti v log2 m krocích) • uloží pouze jedno číslo odpovídající Yl!h=bmXx y> Výhodnější z hlediska paměťových přenosů i spouštění kernelů • čteme „ + i + + .. + -t£_ = (n - 1)-^ mm1 m'°&mn v ' m—l • přibližně n + ^ čtení, ^ zápisů • logm n spuštění kernelu Hledání slabých míst oooooooo Implementace 1 .global__ void reducel(int *v){ extern__shared__ int sv [] ; unsigned int tid = threadldx.x; unsigned int i = blockldx.x*blockDim.x + threadldx.x; sv [ t id ] = v [ i ] ; __syncthreads (); for(unsigned int s=l; s < blockDim.x; s *= if (tid % (2*s) == 0) s v [tid] += s v [ t i d + s]; __syncthreads(); } 2) { if (tid = 0) v[blockldx.x] = sv[0]; í -00.0 Jiří Filipovič Optimalizace Revize násobení matic ooooooooooooooo Výkon Redukce Obecné zásady OOOOOO^OOOOOOOOOOOO OOOOOOO Hledání slabých míst OOOOOOOO Vysoká úroveň divergence • 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.77GB/s, 0.94MEIem/s. 4 □ ► 4 fiP ► 4 m -00,0 Revize násobení matic ooooooooooooooo Implementace 2 Redukce Obecné zásady OOOOOOOÄOOOOOOOOOOO OOOOOOO Hledání slabých míst OOOOOOOO 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.33GB/s, 2.08MEIem/s. Řeší divergenci, generuje konflikty bank. 2) { 4 □ ► 4 fiP ► 4 m -00,0 Revize násobení matic ooooooooooooooo Implementace 3 Redukce Obecné zásady 00000000*0000000000 OOOOOOO Hledání slabých míst OOOOOOOO Tak ještě jinak... for (unsigned int s = blockDim . x/2; s > 0; s »= 1) { if (tid < s) s v [tid] += s v[t i d + s]; __syncthreads () ; } Žádná divergence ani konflikty. Přenos 16.34 GB/s, 4.08MEIem/s. Polovina threadů nic nepočítá... i -00.0 Revize násobení matic ooooooooooooooo Implementace 4 Redukce Obecné zásady OOOOOOOOOÄOOOOOOOOO OOOOOOO Hledání slabých míst OOOOOOOO První sčítání provedeme již během načítání. unsigned int i sv[t id 1 = v [ i I blockldx.x*(blockDim.x * 2) v[i+blockDim.x 1; threadldx.x; Přenos 27.16 GB/s, 6.79MEIem/s. Data zřejmě čteme optimálně, stále je zde však výkonová rezerva zaměřme se na instrukce. i -00.0 Revize násobení matic ooooooooooooooo Implementace 5 Redukce Obecné zásady 0000000000*00000000 OOOOOOO Hledání slabých míst OOOOOOOO V jednotlivých krocích redukce ubývá aktivních threadů • nakonec bude pracovat pouze jeden warp • ten je však synchronizován implicitně, můžeme tedy odebrat syncthreadsQ • podmínka if[tid < s) je zde zbytečná (nic neušetří) Unrollujme tedy poslední warp... Revize násobení matic ooooooooooooooo Implementace 5 for (unsigned int s = blockDim . x/2; s > 32; s »= 1){ if (tid < s) sv [tid] += s v[t i d + s]; __syncthreads () ; } if (tid < 32){ SV tid] += SV tid 4 - 32] SV tid] += SV tid 4 - 16] SV tid] += SV tid 4 - 8]; SV tid] += SV tid 4 - 4]; SV tid] += SV tid 4 - 2]; SV tid] += SV tid 4 - i]; } Ušetříme čas i ostatním waprům (zkončí dříve s for cyklem). Přenos 37.68 GB/s, 9.42MEIem/s. Redukce Obecné zásady OOOOOOOOOOOÄOOOOOOO OOOOOOO Hledání slabých míst OOOOOOOO Jiří Filipovič Optimalizace Revize násobení matic ooooooooooooooo Implementace 6 Jak je to s rozvinutím for cyklu? Známe-li počet iterací, můžeme cyklus rozvinout • počet iterací je závislý na velikosti bloku Můžeme být obecní? • algoritmus pracuje s bloky o velikosti 2" • velikost bloku je shora omezena • známe-li při kompilaci velikost bloku, můžeme použít šablonu template __global__ void reduce6(int *v) Jiří Filipovič Optimalizace Redukce Obecné zásady 000000000000*000000 OOOOOOO Hledání slabých míst OOOOOOOO Revize násobení matic ooooooooooooooo Implementace 6 Podmínky s blockSize se vyhodnotí již pří překladu: if (blockSize >= 512){ if (tid < 256) sv[tid] += sv[tid + 256]; __syncthreads () ; } if (blockSize >= 256){ if (tid < 128) sv[tid] += sv[tid + 128]; __syncthreads () ; } if (blockSize >= 128){ if (tid < 64) sv[tid] += sv[tid + 64]; __syncthreads () ; } Jiří Filipovič Optimalizace Redukce Obecné zásady OOOOOOOOOOOOO0OOOOO OOOOOOO Hledání slabých míst OOOOOOOO Revize násobení matic ooooooooooooooo Implementace 6 Redukce Obecné zásady OOOOOOOOOOOOOOÄOOOO OOOOOOO Hledání slabých míst OOOOOOOO if (tid < 32){ if (blocksize >= 64) s v[tid] += s v[tid 4- 32] if (blocksize >= 32) s v[tid] += s v[tid 4- 16] if (blocksize >= 16) s v[tid] += s v[tid 4- 8]; if (blocksize >= 8) s v[tid] += sv[tid + 4]; if (blocksize >= 4) s v[tid] 4-= sv[tid + 2]; if (blocksize >= 2) s v[tid] 4-= sv[tid + i]; } Spustení kernelu: reduce6»(d_v ) ; Přenos 50.64 GB/s, 12.66 MEIem/s. Jiří Filipovič Optimalizace Revize násobení matic Redukce Obecné zásady Hledání slabých míst ooooooooooooooo ooooooooooooooo»ooo ooooooo oooooooo Implementace 7 Můžeme algoritmus ještě vylepšit? Vratme se zpět ke složitosti: • celkem log n kroků • celkem n — 1 sčítání • časová složitost pro p threadů běžících paralelně (p procesorů) 0(* + \ogn) Cena paralelního výpočtu • definována jako počet procesorů krát časová složitost • přidělíme-li každému datovému elementu jeden thread, lze uvažovat p = n • pak je cena 0(n ■ log n) • není efektivní Jiří Filipovič Optimalizace Revize násobení matic Redukce Obecné zásady Hledání slabých míst ooooooooooooooo oooooooooooooooo«oo ooooooo oooooooo Implementace 7 Snížení ceny • použijeme threadů • každý thread provede O(logn) sekvenčních kroků • následně se provede O(logn) 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 • to přináší výhodu v momentě, kdy máme výrazně více threadů, než je třeba k saturaci GPU • navíc snižujeme overhead spouštění kernelů Jiří Filipovič Optimalizace Revize násobení matic ooooooooooooooo Implementace 7 Modifikujeme načítání do sdílené paměti unsigned int gridSize = blockSize*2*gridDim.x; sv[tid] = 0; while(i < n){ sv[tid] += v[i] + v[i+blockSize ] ; i += gridSize; } __syncthreads (); Přenos 77.21 GB/s, 19.3MEIem/s. Jiří Filipovič Optimalizace Redukce Obecné zásady OOOOOOOOOOOOOOOOO^O OOOOOOO Hledání slabých míst OOOOOOOO Revize násobení matic ooooooooooooooo Implementace 7 Modifikujeme načítání do sdílené paměti unsigned int gridSize = blockSize*2*gridDim.x; sv[tid] = 0; while(i < n){ sv[tid] += v[i] + v[i+blockSize ] ; i += gridSize; } __syncthreads (); Přenos 77.21 GB/s, 19.3MEIem/s. Jednotlivé implementace jsou k nalezení v CUDA SDK. Jiří Filipovič Optimalizace Redukce Obecné zásady 00000000000000000*0 OOOOOOO Hledání slabých míst OOOOOOOO Revize násobení matic ooooooooooooooo Redukce Obecné zásady OOOOOOOOOOOOOOOOOO* ooooooo Hledání slabých míst oooooooo Poznámky pro c.c. 2.0 Kompilátor může odkládat uložení dat do sdílené paměti • při unrollingu posledního warpu je zapotřebí použít volatile proměnnou • je vhodné naznačit kompilátoru, co si smí držet lokálně if (tid < 32){ volatile float *s = sv; if (blockSize >= 64) s [tid — mySum — mySum + s[tid ■f 32]; if (blockSize >= 32) s [tid — mySum — mySum + s[tid ■f 16]; if (blockSize >= 16) s [tid — mySum — mySum + s j tid H- 8]; if (blockSize >= 8) s [tid] — mySum = = mySum + s[t id + 4]; if (blockSize >= 4) s [tid] — mySum = = mySum + s j t id + 2]; if (blockSize >= 2) s [tid] — mySum = = mySum + s[t id + i]; } Jiří Filipovič Optimalizace Revize násobení matic ooooooooooooooo Redukce Obecné zásady OOOOOOOOOOOOOOOOOOO «000000 Hledání slabých míst oooooooo Výběr vhodného problému Než se pustíme do GPU akcelerace, je vhodné se zamyslet, jestli nám může pomoci :-). Akcelerovaný problém by měl být • kritický pro výkon aplikace • musí se jednat o dostatečně velký problém (z hlediska počtu operací k jeho vyřešení) • musí být paralelizovatelný (to zpravidla velké problémy jsou) • k řešení problému musí být zapotřebí dostatek operací na jeden datový element (omezení přenosu zadání po PCI-E) Optimalizujeme čas či spotřebu? Jiří Filipovič Optimalizace Revize násobení matic Redukce Obecné zásady Hledání slabých míst OOOOOOOOOOOOOOO OOOOOOOOOOOOOOOOOOO 0*00000 oooooooo Návrh algoritmu Paralelizace • v řešeném problému je třeba najít paralelismus • již zde je vhodné uvažovat o omezeních architektury Obtížně akcelerovatelné jsou algoritmy, pro které platí • jednotlivé thready přistupují na náhodná místa paměti • silně divergentní běh kódu • nedostatečný paralelismus či složitá synchronizace Redukce Obecné zásady OOOOOOOOOOOOOOOOOOO OOÍOOOO Revize násobení matic ooooooooooooooo Prevence chyb během implementace Hledání slabých míst OOOOOOOO Testujte úspěšnost volání API a kernelů • chyby se jinak umí projevit se spozdením Na GPU je docela deterministická alokace paměti • pokud se nic nezapíše, obvykle dostanete výsledek z minulého běhu algoritmu • pro účely ladění mažte výstupní data Pozor na pošlapání sdílené paměti • kernel často nespadne, objevují se interference mezi bloky Jiří Filipovič Optimalizace Revize násobení matic Redukce Obecné zásady Hledání slabých míst OOOOOOOOOOOOOOO OOOOOOOOOOOOOOOOOOO OOOÄOOO oooooooo Optimalizace Je ruzumné postupovat od obecně významnějších k méně významným (tak se jejich efekt lépe projeví) • přístup do globální paměti (bandwidth, latence) • přístup do ostatních pamětí • konfigurace běhu (počet threadů na blok, množství práce na thread) • divergence běhu • optimalizace na úrovni instrukcí Je dobré psát snadno konfigurovatelný kód • konfigurovatelná velikost bloku, práce na thread, .. . • konfigurace výhodnější přes makra (část výpočtů proveditelná v době kompilace) Jiří Filipovič Optimalizace Revize násobení matic Redukce Obecné zásady Hledání slabých míst OOOOOOOOOOOOOOO OOOOOOOOOOOOOOOOOOO OOOOÍOO oooooooo Pozor na interpretaci rychlosti algoritmu Efekt některých optimalizací může být skryt významnějšími neoptimalitami • omezíme přednostním aplikováním významnějších optimalizací • omezíme používáním profileru Prostor optimalizací je nespojitý • dáno omezenými zdroji GPU » rychlejší kód threadu může vézt k celkově nižšímu výkonu Výkon je závislý na velikosti problému • nedostatečné pokrytí multiprocesorů • partition camping Revize násobení matic Redukce ooooooooooooooo ooooooooooooooooooo Obecné zásady ooooo«o Hledání slabých míst OOOOOOOO Jaké zrychlení oproti CPU je reál né? Základní odhad zrychlení vychází z porovnání aritmetického výkonu a propustnosti paměti • GPU však nemusí přinést adekvátní zrychlení • nedostatečně či nevhodně paralelizovatelný algoritmus • nevhodné datové struktury, náhodný přístup • PCI-E bottleneck (málo výpočtu vzhledem k přenosům, multi-GPU algoritmy) • GPU také může přinést vyšší zrychlení • významné využití SFU • komplikovaná vektorizace u CPU • degradace výkonu paměti u CPU • špatně škálující SMP • odlišné škálování CPU a GPU s rostoucí velikostí problému Jiří Filipovič Optimalizace Redukce Obecné zásady OOOOOOOOOOOOOOOOOOO 000000« Revize násobení matic ooooooooooooooo Pozor na příliš optimistická měření Hledání slabých míst OOOOOOOO Vysoké zrychlení má často za příčinu špatný CPU algoritmus • je zapotřebí si uvědomit, že CPU má více jader a vektorové jednotky • nevektorizovaný jednothreadový kód využívá (v jednoduché přesnosti) 1/16 teoretického maxima 4-jádrového CPU a 1/64 8-jádrového s AVX instrukcema • přínos GPU řešení lze podložit spočítáním flopsů Jiří Filipovič Optimalizace Hledání slabých míst •ooooooo Ladění výkonu Základní výpočet aritmetických operácia paměťových přenosů nám říká, kde jsou maxima algoritmu a na co se primárně zaměřit • někdy není jasný bottleneck konkrétní implementace (zpravidla máme více instrukcí, než aritmetických operací nutných pro řešení problému) • profiling kódu - vhodný pro identifikaci problémů s propustností instrukcí či paměti, slabý pro identifikaci problémů s latencí • modifikace kódu - přesnější, ale náročnější metoda, není použitelná vždy i -00.0 Revize násobení matic Redukce Obecné zásady Hledání slabých míst ooooooooooooooo ooooooooooooooooooo ooooooo o«oooooo Profiling Jak blízko jsme k maximu HW? • IPC - pro Fermi, počet instrukcí na cyklus, maximum 2 (nebereme-li v úvahu mix instrukcí) • instruction throughput - pro c.c. 1.x, procento maximální rychlosti spouštění (single-issue) instrukci • pro Fermi jsou také reportovány rychlosti přenosu paměti (pro c.c. 1.x nutno dopočítat) Revize násobení matic ooooooooooooooo Redukce Obecné zásady ooooooooooooooooooo ooooooo Hledání slabých míst ooaooooo Profiling Získání aktuálního poměru mezi instrukcemi a paměťovými přenosy: • instructionJssued udává počet emitovaných instrukcí na warp na multiprocesor • dram^reads a dram-writes udává počet 32-bytových přenosů • poměr instrukcí k paměťovým přenosům získáme pomocí #SM-32-instruction_issued vzorce- —- 32(dram_reads+dram_writes) • lze použít také instruction/byte, používá analogický výpočet, ale pro cacheovaný přístup (je otázka, co chceme změřit) Jiří Filipovič Optimalizace Revize násobení matic Redukce Obecné zásady Hledání slabých míst OOOOOOOOOOOOOOO OOOOOOOOOOOOOOOOOOO OOOOOOO 000*0000 Profiling Serializace • Replayed Instructions - procento instrukcí, které byly vícekrát zavedeny (především způsobeno serializací) • Divergent Branches - procento větvení, které divergovalo • Control Flow Divergence - procento instrukcí, které nebyly prováděny všemi thready ve warpu • Shared Bank Conflict Replay - procento instrukcí znovuzavedených kvůli konfliktu bank sdílené paměti • Shared Memory Bank Conflict per Shared Memory Instruction - procento přístupů do sdílené paměti, které vyvolaly konflikt bank Jiří Filipovič Optimalizace Revize násobení matic Redukce Obecné zásady Hledání slabých míst OOOOOOOOOOOOOOO OOOOOOOOOOOOOOOOOOO OOOOOOO 0000*000 Profiling Přístup do paměti • Global memory excess had - pro Fermi, procento nadbytečně přenášených dat, analogicky pro store • * hit ratio - procento přístupů realizovaných přes příslušnou cache Mnoho dalších užitečných ukazatelů - viz manuál. Jiří Filipovič Optimalizace Revize násobení matic Redukce Obecné zásady Hledání slabých míst OOOOOOOOOOOOOOO OOOOOOOOOOOOOOOOOOO OOOOOOO 00000*00 Modifikace kódu Výkon paměťových přenosů • zakomentujeme výpočet • načtená data musíme nějak "použít" • kontrola profilerem, že přenášíme stále stejné množství dat Výkon samostatného výpočtu • odstraníme přenosy dat • výsledek výpočtu je však třeba uložit, aby kompilátor neodstranil výpočet • my ale ukládat nechceme.. . • uložení výsledků lze vložit do podmínky, která nikdy nebude splněna • u rychlých kernelů pozor na overhead jejich spuštění Jiří Filipovič Optimalizace Revize násobení matic Redukce Obecné zásady Hledání slabých míst ooooooooooooooo ooooooooooooooooooo ooooooo ooooooao Modifikace kódu Pozor na změnu dostupného paralelismu • pokud modifikace kódu ubere využití zdrojů GPU • můžeme omezit paralelismus přidáním dynamicky alokované sdílené paměti ke spuštění kernelu Interpretace naměřených rychlostí • celkový čas běhu se blíží součtu času výpočtu a přenosů paměti - problém s latencí • jeden z časů převládá a blíží se celkovému času běhu -rychlost je omezena výpočtem nebo pamětí, víme kam zaměřit optimalizaci • oba časy podobné a blízké celkovému času běhu - jediná možnost pro zrychlení je optimalizovat obojí Jiří Filipovič Optimalizace Revize násobení matic Redukce Obecné zásady Hledání slabých míst OOOOOOOOOOOOOOO OOOOOOOOOOOOOOOOOOO OOOOOOO OOOOOOO* Modifikace kódu Odhad dopadu optimalizace • máme-li identifikován výkonnostní problém • chceme odhadnout dopad optimalizace, než se do ní pustíme • „zmrzačení kódu" - úprava, porušující korektnost, ale odstraňující neoptimalitu • nelze vždy • může rychle ukázat, že cílíme na nesprávný výkonnostní problém