Revize násobení matic Redukce Obecné zásady Pár praktických rad ooooooooooooooo oooooooooooooooooo oooooo ooo Optimalizace Jiří Filipovič podzim 2010 Jiří Filipovič Optimalizace Revize násobení matic Redukce Obecné zásady Pár praktických rad •oooooooooooooo oooooooooooooooooo oooooo ooo Naivní implementace 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++) t mp += A[y * n+k ] * B[k * n+x ]; C[y*n + x] = tmp; Jiří Filipovič Optimalizace Revize násobení matic Redukce Obecné zásady Pár praktických rad o»ooooooooooooo oooooooooooooooooo oooooo ooo 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 Jiří Filipovič Optimalizace □ gP - = Revize násobení matic Redukce Obecné zásady Pár praktických rad o»ooooooooooooo oooooooooooooooooo oooooo ooo 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 neprekladaný, v opačném případě je prokládaný □ - = = ^q^o Jiří Filipovič Optimalizace Revize násobení matic oo»oooooooooooo Co jsme se naučili Redukce Obecné zásady oooooooooooooooooo oooooo Pár praktických rad ooo 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... Jiří Filipovič Optimalizace □ gP - = Revize násobení matic Redukce Obecné zásady Pár praktických rad ooo»ooooooooooo oooooooooooooooooo oooooo ooo 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 = for (int b = As[ty][tx] Bs[ty][tx] _syncthreads(); O.Of ; 0; b < n/BLOCK ; b++){ = A[(ty + by*BL0CK)*n = B[(ty + b*BL0CK)*n 4 I- b*BL0CK+tx ] ; bx*BL0CK+tx] ; } for (int k = 0; k < BLOCK; k++) Csub += As[ty][k]* Bs[k][tx]; __syncthreads(); C[(ty + by*BL0CK)*n + bx*BL0CK+tx] Csub ; □ S Jiří Filipovič Optimalizace Revize násobení matic Redukce Obecné zásady Pár praktických rad oooo»oooooooooo oooooooooooooooooo oooooo ooo Hříchy implementace As[ty][tx] = A[(ty + by*BL0CK)*n + b*BL0CK+tx]; Bs[ty][tx] = 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 Jiří Filipovič Optimalizace Redukce Obecné zásady oooooooooooooooooo oooooo Revize násobení matic ooooo»ooooooooo Teoretické maximum Pár praktických rad ooo 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 415 GFIops • stále jsme od něj však daleko Jiří Filipovič Optimalizace □ gP - = Revize násobení matic Redukce Obecné zásady Pár praktických rad oooooo»oooooooo oooooooooooooooooo oooooo ooo 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 □ - = = ^q^o Jiří Filipovič Optimalizace Revize násobení matic ooooooo»ooooooo Redukce Obecné zásady oooooooooooooooooo oooooo Pár praktických rad ooo Nalezení lepší implementace Lze počet load instrukcí omezit? Jiří Filipovič Optimalizace Revize násobení matic ooooooo»ooooooo Redukce Obecné zásady oooooooooooooooooo oooooo Pár praktických rad ooo Nalezení lepší implementace Lze počet load instrukcí omezit? • data ve sdílené paměti snižují přenosy z paměti globální Jiří Filipovič Optimalizace □ gP - = Revize násobení matic ooooooo»ooooooo Redukce Obecné zásady oooooooooooooooooo oooooo Pár praktických rad ooo 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? Jiří Filipovič Optimalizace □ gP - = Revize násobení matic ooooooo»ooooooo Redukce Obecné zásady oooooooooooooooooo oooooo Pár praktických rad ooo 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 □ gP - = Revize násobení matic ooooooo»ooooooo Redukce Obecné zásady oooooooooooooooooo oooooo Pár praktických rad ooo 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 overheadu daném pointerovou aritmetikou • experimentálně najdeme vhodnou velikost bloku □ - = = ^q^o Jiří Filipovič Optimalizace Revize násobení matic Redukce Obecné zásady Pár praktických rad oooooooo»oooooo oooooooooooooooooo oooooo ooo Nalezení lepší implementace Nejlepší výsledky dosaženy pro bloky velikosti 32 x 32, na kterých pracuje 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ě Jiří Filipovič Optimalizace □ gP - = Revize násobení matic Redukce Obecné zásady Pár praktických rad ooooooooo«ooooo oooooooooooooooooo oooooo ooo 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[$ofsl+0x0008], $r0, $r7 mad.rn.f32 $r8, s[$ofs3+0x0008], $r0, $r8 Jiří Filipovič Optimalizace Revize násobení matic ooooooooo«ooooo Redukce Obecné zásady oooooooooooooooooo oooooo Pár praktických rad ooo 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[$ofsl+0x0008], $r0, $r7 mad.rn.f32 $r8, s[$ofs3+0x0008], $r0, $r8 Kompilátor dokázal převést adresaci přes k na konstantní offsety pouze u proměnné As • k Bs je přistupováno prokládané • znamená to jednu add instrukci navíc □ - = = ^q^o Jiří Filipovič Optimalizace Revize násobení matic oooooooooo»oooo Redukce Obecné zásady oooooooooooooooooo oooooo Pár praktických rad ooo 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 ] ; Jiří Filipovič Optimalizace Revize násobení matic oooooooooo»oooo Redukce Obecné zásady oooooooooooooooooo oooooo Pár praktických rad ooo 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[$ofsl+0x0008], $r0, $r8 Jiří Filipovič Optimalizace Revize násobení matic oooooooooo»oooo Redukce Obecné zásady oooooooooooooooooo oooooo Pár praktických rad ooo 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[$ofsl+0x0008], $r0, $r8 Nový problém - konflikty bank sdílené paměti Jiří Filipovič Optimalizace Revize násobení matic oooooooooo»oooo Redukce Obecné zásady oooooooooooooooooo oooooo Pár praktických rad ooo 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[$ofsl+0x0008], $r0, $r8 Nový problém - konflikty bank sdílené paměti • vyřeší padding Jiří Filipovič Optimalizace Revize násobení matic Redukce Obecné zásady Pár praktických rad oooooooooo»oooo oooooooooooooooooo oooooo ooo 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[$ofsl+0x0008], $r0, $r8 Nový problém - konflikty bank sdílené paměti • vyřeší padding Výsledná rychlost: 276.2 GFIops. Jiří Filipovič Optimalizace Revize násobení matic Redukce Obecné zásady Pár praktických rad ooooooooooo«ooo oooooooooooooooooo oooooo ooo Lze matice násobit ještě rychleji? 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 □ gP - = Revize násobení matic ooooooooooo«ooo Redukce Obecné zásady oooooooooooooooooo oooooo Lze matice násobit ještě rychleji? Pár praktických rad ooo 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 □ gP - = Revize násobení matic ooooooooooo«ooo Redukce Obecné zásady oooooooooooooooooo oooooo Lze matice násobit ještě rychleji? Pár praktických rad ooo 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 □ gP - = Revize násobení matic Redukce Obecné zásady Pár praktických rad oooooooooooo»oo oooooooooooooooooo oooooo ooo Přehodnocený blokový přístup Namísto čtvercových bloků v matici C můžeme použít obdélníkové Jiří Filipovič Optimalizace □ r5P - = Revize násobení matic oooooooooooo»oo Redukce Obecné zásady oooooooooooooooooo oooooo Pár praktických rad ooo 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 Jiří Filipovič Optimalizace □ gP - = Revize násobení matic oooooooooooo»oo Redukce Obecné zásady oooooooooooooooooo oooooo Pár praktických rad ooo 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 Jiří Filipovič Optimalizace □ gP - = Revize násobení matic oooooooooooo»oo Redukce Obecné zásady oooooooooooooooooo oooooo Pár praktických rad ooo 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 Jiří Filipovič Optimalizace □ gP - = Revize násobení matic oooooooooooo»oo Redukce Obecné zásady oooooooooooooooooo oooooo Pár praktických rad ooo 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 Jiří Filipovič Optimalizace □ gP - = Revize násobení matic oooooooooooo»oo Redukce Obecné zásady oooooooooooooooooo oooooo Pár praktických rad ooo 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 Jiří Filipovič Optimalizace □ gP - = Revize násobení matic oooooooooooo»oo Redukce Obecné zásady oooooooooooooooooo oooooo Pár praktických rad ooo 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 a není nutná aritmetika uprostřed smyčky (viz předchozí optimalizace) Jiří Filipovič Optimalizace □ r5P - = Revize násobení matic oooooooooooo»oo Redukce Obecné zásady oooooooooooooooooo oooooo Pár praktických rad ooo 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 a 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 □ - = = ^q^o Jiří Filipovič Optimalizace Revize násobení matic ooooooooooooo»o Implementace Redukce Obecné zásady oooooooooooooooooo oooooo Pár praktických rad ooo 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 □ gP - = Revize násobení matic ooooooooooooo»o Implementace Redukce Obecné zásady oooooooooooooooooo oooooo Pár praktických rad ooo 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 □ gP - = Revize násobení matic 0OOOOOOOOOOOOO» Redukce Obecné zásady oooooooooooooooooo oooooo Pár praktických rad ooo Shrnutí Implementace rychlost rel. A abs. A Naivní implementace, thready 1 x 128 3.9GFIops 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 71 x Jen jeden blok ve sdílené paměti 375 GFIops 1.36x 96 x Jiří Filipovič Optimalizace Revize násobení matic 0OOOOOOOOOOOOO» Redukce Obecné zásady oooooooooooooooooo oooooo Pár praktických rad ooo Shrnutí Implementace rychlost rel. A abs. A Naivní implementace, thready 1 x 128 3.9GFIops 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 71 x 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 □ gP - = Revize násobení matic 0OOOOOOOOOOOOO» Redukce Obecné zásady oooooooooooooooooo oooooo Pár praktických rad ooo Shrnutí Implementace rychlost rel. A abs. A Naivní implementace, thready 1 x 128 3.9GFIops 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 71 x 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í. □ - = = ^q^o Jiří Filipovič Optimalizace Revize násobení matic Redukce Obecné zásady Pár praktických rad OOOOOOOOOOOOOOO »00000000000000000 oooooo ooo Součet prvků vektoru Pro vektor von prvcích chceme spočítat x = ^"=1 vi- Jiří Filipovič Optimalizace □ gP - = Revize násobení matic Redukce Obecné zásady Pár praktických rad OOOOOOOOOOOOOOO »00000000000000000 oooooo ooo Součet prvků vektoru Pro vektor v o n prvcích chceme spočítat x = J^/Li vi-Zápis v jazyce C int x = O; for (int i = 0; i < n; i++) x += v[i] ; Jednotlivé iterace cyklu jsou na sobě závislé. Jiří Filipovič Optimalizace □ gP - = Revize násobení matic ooooooooooooooo Redukce Obecné zásady •ooooooooooooooooo oooooo Součet prvků vektoru Pár praktických rad ooo Pro vektor v o n prvcích chceme spočítat x = J^/Li vi-Zápis v jazyce C int x = 0; for (int i = 0; i < n; i++) x += v[i] ; Jednotlivé iterace cyklu jsou na sobě závislé. • nemůžeme udělat všechnu práci paralelně » sčítání je však (alespoň teoreticky :-)) asocitativní • není tedy nutno počítat sekvenčně Jiří Filipovič Optimalizace □ gP - = Revize násobení matic Redukce Obecné zásady Pár praktických rad ooooooooooooooo o»oooooooooooooooo oooooo ooo Paralelní algoritmus Představený sekvenční algoritmus provádí pro 8 prvků výpočet: (((((( ví + ^2) + v3) + 1/4) + v5) + v6) + vj) + v8 Jiří Filipovič Optimalizace □ gP - = Revize násobení matic Redukce Obecné zásady Pár praktických rad ooooooooooooooo o»oooooooooooooooo oooooo ooo Paralelní algoritmus Představený sekvenční algoritmus provádí pro 8 prvků výpočet: (((((( ví + ^2) + v3) + 1/4) + v5) + v6) + vj) + v8 Sčítání je asociativní... spřeházejme tedy závorky: ((1/1 + v2) + (1/3 + vA)) + ((1/5 + v6) + {vj + v8)) Jiří Filipovič Optimalizace □ r5P - = Revize násobení matic ooooooooooooooo Redukce Obecné zásady o»oooooooooooooooo oooooo Pár praktických rad ooo Paralelní algoritmus Představený sekvenční algoritmus provádí pro 8 prvků výpočet: (((((( ví + v2) + v3) + 1/4) + v5) + v6) + vj) + v8 Sčítání je asociativní... spřeházejme tedy závorky: ((1/1 + v2) + (1/3 + vA)) + ((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 ooooooooooooooo Paralelní algoritmus Redukce Obecné zásady oo»ooooooooooooooo oooooo Pár praktických rad ooo 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 □ r5P - = Revize násobení matic ooooooooooooooo Redukce Obecné zásady ooo»oooooooooooooo oooooo Pár praktických rad ooo Naivní přístup Nejjednodušší schéma algoritmu: • kernel pro sudá / < 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. Jiří Filipovič Optimalizace □ gP - = Revize násobení matic Redukce Obecné zásady Pár praktických rad ooooooooooooooo oooo«ooooooooooooo oooooo ooo 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í YlT^mXx vi Výhodnější z hlediska paměťových přenosů i spouštění kernelů • přibližně n + čtení, ^ zápisů • logm n spuštění kernelu Jiří Filipovič Optimalizace □ gP - = Revize násobení matic Redukce Obecné zásady ooooooooooooooo ooooo»oooooooooooo oooooo Pár praktických rad ooo Implementace 1 __global__ void reducel(int *v){ extern __shared__ int sv[]; unsigned int tid = threadldx.x; unsigned int i = blockldx.x*blockDim.x + sv[t id ] = v[i ] ; __syncthreads(); threadldx.x; for(unsigned int s=l; s < blockDim.x; s * if (tid % (2*s) == 0) sv[tid] += sv[tid + s]; syncthreads(); } = 2) { if (tid = 0) v[blockldx.x] = sv[0]; } □ gp - = = ^)c^o Jiří Filipovič Optimalizace Revize násobení matic Redukce Obecné zásady Pár praktických rad ooooooooooooooo oooooo«ooooooooooo oooooo ooo Výkon 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.77 GB/s, 0.94 MElem/s. Jiří Filipovič Optimalizace □ gP - = Revize násobení matic ooooooooooooooo Redukce Obecné zásady ooooooo»oooooooooo oooooo Pár praktických rad ooo Implementace 2 Nahradíme indexaci ve for cyklu for (unsigned int s = 1; s < blockDim.x; s *= 2) { int index = 2 * s * tid; if (index < blockDim.x) sv[index] += sv[index + s]; __syncthreads (); } Přenos 8.33 GB/s, 2.08 MEIem/s. Řeší divergenci, generuje konflikty bank. Jiří Filipovič Optimalizace Revize násobení matic ooooooooooooooo Implementace 3 Redukce Obecné zásady oooooooo»ooooooooo oooooo Pár praktických rad ooo Tak ještě jinak... for (unsigned int s = blockDim. x/2; s > 0; s »= 1) { if (tid < s) sv[tid] += sv[tid + s]; __ sync t hr e ads (); } Žádná divergence ani konflikty. Přenos 16.34 GB/s, 4.08 MElem/s. Polovina threadů nic nepočítá... Jiří Filipovič Optimalizace Revize násobení matic Redukce Obecné zásady Pár praktických rad ooooooooooooooo ooooooooo»oooooooo oooooo ooo Implementace 4 První sčítání provedeme již během načítání. unsigned int i = blockldx.x*(blockDim.x*2) + threadldx.x; sv[tid] = v[i] + v[i+blockDim.x ] ; Přenos 27.16 GB/s, 6.79 MElem/s. Data zřejmě čteme optimálně, stále je zde však výkonová rezerva -zaměřme se na instrukce. Jiří Filipovič Optimalizace □ r5P - = Revize násobení matic ooooooooooooooo Implementace 5 Redukce Obecné zásady oooooooooosooooooo oooooo Pár praktických rad ooo 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... Jiří Filipovič Optimalizace □ gP - = Revize násobení matic ooooooooooooooo Redukce Obecné zásady ooooooooooo»oooooo oooooo Pár praktických rad ooo Implementace 5 for (unsigned int s = blockDim. x/2; s > 32; s »= 1){ if (tid < s) sv[tid] += sv[tid + s]; __syncthreads (); } if (tid < 32){ SV tid] += SV tid ^ - 32] SV tid] += SV tid ^ - 16] SV tid] += SV tid ^ - 8]; SV tid] += SV tid ^ " 4]; SV tid] += SV tid ^ " 2]; SV tid] += SV tid ^ - i]; } Ušetříme čas i ostatním waprům (zkončí dříve s for cyklem). Přenos 37.68 GB/s, 9.42 MElem/s. Jiří Filipovič Optimalizace Revize násobení matic ooooooooooooooo Redukce Obecné zásady oooooooooooosooooo oooooo Pár praktických rad ooo 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 □ gP - = Revize násobení matic ooooooooooooooo Redukce Obecné zásady ooooooooooooo»oooo oooooo Pár praktických rad ooo 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 Revize násobení matic ooooooooooooooo Implementace 6 Redukce Obecné zásady oooooooooooooo»ooo oooooo Pár praktických rad ooo if (tid < 32){ if (blockSize >= 64) sv[t id] += sv[t id 4- 32] if (blockSize >= 32) sv[t id] += sv[t id 4- 16] if (blockSize >= 16) sv[t id] += sv[t id 4- 8]; if (blockSize >= 8) sv[t id] ■f= sv [ tid + 4]; if (blockSize >= 4) sv[t id] ■f= sv [ tid + 2]; if (blockSize >= 2) sv[t id] ■f= sv [ tid + l]: } Spuštění kernelu: reduce6»(d_v ); Přenos 50.64GB/s, 12.66MElem/s. Jiří Filipovič Optimalizace Revize násobení matic ooooooooooooooo Redukce Obecné zásady ooooooooooooooo»oo oooooo Pár praktických rad ooo Implementace 7 Můžeme algoritmus ještě vylepšit? Vratme se zpět ke složitosti: • celkem logn kroků • celkem n — 1 sčítání • časová složitost pro p threadů běžících paralelně (p procesorů) 0(£ + logn) 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 rí) • není efektivní □ g - = = ^q^O Jiří Filipovič Optimalizace Revize násobení matic Redukce Obecné zásady Pár praktických rad ooooooooooooooo oooooooooooooooo»o oooooo ooo Implementace 7 Snížení ceny • použijeme O(j^) threadů • každý thread provede (D(logn) sekvenčních kroků • následně se provede (D(logn) paralelních kroků • časová složitost zůstane • cena se sníží na 0{rí) 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ů □ - = = ^q^o Jiří Filipovič Optimalizace Revize násobení matic ooooooooooooooo Redukce Obecné zásady 00000000000000000» oooooo Pár praktických rad ooo 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.3 MElem/s. Jiří Filipovič Optimalizace Revize násobení matic ooooooooooooooo Redukce Obecné zásady 00000000000000000» oooooo Pár praktických rad ooo 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.3 MElem/s. Jednotlivé implementace jsou k nalezení v CUDA SDK. Jiří Filipovič Optimalizace Revize násobení matic ooooooooooooooo Redukce Obecné zásady OOOOOOOOOOOOOOOOOO »00000 Pár praktických rad ooo 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 □ gP - = Revize násobení matic ooooooooooooooo Redukce Obecné zásady oooooooooooooooooo o»oooo Pár praktických rad ooo Postup návrhu algoritmu Paralelizace • v řešeném problému je třeba najít paralelismus • již zde je vhodné uvažovat o omezeních architektury Teoretické maximum rychlosti algoritmu • než začneme implementovat, je vhodné mít představu, jak rychle může algoritmus na daném HW pracovat • základní omezení dává paměťová propustnost a aritmetický výkon • výkon CPU a GPU se může sčítat Jiří Filipovič Optimalizace □ gP - = Revize násobení matic Redukce Obecné zásady Pár praktických rad ooooooooooooooo oooooooooooooooooo oo»ooo ooo 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í Nezapomínejte • poctivě benchmarkovat • kontrolovat kód profilerem □ - = = ^q^o Jiří Filipovič Optimalizace Revize násobení matic ooooooooooooooo Redukce Obecné zásady oooooooooooooooooo ooo»oo Pár praktických rad ooo 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 • menší instance mohou mít jiné nároky • partition camping Jiří Filipovič Optimalizace □ gP - = Revize násobení matic ooooooooooooooo Redukce Obecné zásady oooooooooooooooooo oooo«o Pár praktických rad ooo Jaké zrychlení oproti CPU je reálné? • 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 a odlišné škálování CPU a GPU s rostoucí velikostí problému □ - = = ^q^o Jiří Filipovič Optimalizace Revize násobení matic ooooooooooooooo Redukce Obecné zásady OOOOOOOOOOOOOOOOOO 00000» Pozor na příliš optimistická měření Pár praktických rad ooo • 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 • přínos GPU řešení lze podložit spočítáním flopsů Jiří Filipovič Optimalizace □ gP - = Revize násobení matic ooooooooooooooo Prevence chyb Redukce Obecné zásady oooooooooooooooooo oooooo Pár praktických rad •oo • 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 správný výsledek • 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 □ r5P - = Revize násobení matic ooooooooooooooo Redukce Obecné zásady oooooooooooooooooo oooooo Pár praktických rad o«o Ladění výkonu • 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) • někdy není jasný bottleneck a přístup do paměti nemusí být zcela optimální, rychlost instrukcí obtížně odhadnutelná, schopnost překrýt výpočet a paměťové přenosy také • a kdo tedy zdržuje? • paměťové přenosy i výpočet lze benchmarkovat zvlášť • pak můžeme kernel rozdělit, zvětšit, zvýšit obsazenost GPU... Jiří Filipovič Optimalizace □ gP - = Revize násobení matic ooooooooooooooo Redukce Obecné zásady oooooooooooooooooo oooooo Pár praktických rad oo» Měření výkonu • samostatné paměťové přenosy • za komentujem e výpočet • načtená data musíme nějak " použít" (často jednodušší situace u sdílené paměťi) • samostatný výpočet • data nemusíme načítat • 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