GPU hardware Paralelismus Paměťová hierarchie Synchronizace Násobení matic 000 OOOOOOO OOOOOOOOOO OOOOOOOO OOOOOOOOOOO GPU hardware a paralelismus Jiří Filipovič podzim 2013 Jiří Filipovič GPU hardware a paralelismus GPU hardware •OO Paralelismus Paměťová hierarchie OOOOOOO OOOOOOOOOO Synchronizace OOOOOOOO Násobení matic OOOOOOOOOOO Rozdíly v jednotlivých CUDA-GPU Nové generace přinášejí vyšší výkon a výpočetní schopnosti. • výpočetní schopnosti compute capability představují bohatost instrukční sady GPU a množství zdrojů (registry, současně běžící thready aj.) • výkon roste se schopností umístit na jedno GPU více jader V rámci generace se významně liší výkon. • kvůli nabídce levnějších variant • díky pozdějším změnám výrobního procesu • kvůli spotřebě u mobilních GPU Jiří Filipovič GPU hardware a paralelismus GPU hardware Paralelismus Paměťová hierarchie Synchronizace Násobení matic 0«0 OOOOOOO OOOOOOOOOO OOOOOOOO OOOOOOOOOOO Dnes dostupné GPU Dnes dostupné GPU • compute capability 1.0 - 3.5 • s rozdíly se budeme postupně seznamovat • 1 - 30 multiprocesorů (19.2 - 4 500GFLOPs) • frekvence 800 MHz - 1.836 GHz • šířka a rychlost datové sběrnice (64bit - 512bit, 6.4 -288.4 GB/s) Jiří Filipovič GPU hardware a paralelismus GPU hardware Paralelismus Paměťová hierarchie Synchronizace Násobení matic OO* OOOOOOO OOOOOOOOOO OOOOOOOO OOOOOOOOOOO Dostupná řešení Grafické karty GeForce • mainstreamové řešení pro hráče • levné, široce rozšířené, široké spektrum výkonu • nevýhoda - různá omezení (RAM, DP, GPUDirect) Profesionální karty Quadro • z hlediska CUDA stejné jako GeForce • věřší paměť • násobně dražší Tesla • řešení speciálně pro výpočty v CUDA • vždy velká paměť, často vyšší DP výkon aj. • k dispozici jako karty do PCI-E, nebo jako samostatné multi-GPU stroje • taktéž drahé, vhodné pro výpočetní centra či osobní superpočítače i -00.0 Jiří Filipovič GPU hardware a paralelismus GPU hardware Paralelismus Paměťová hierarchie Synchronizace Násobení matic #000000 OOOOOOOOOO OOOOOOOO OOOOOOOOOOO Paralelismus GPU Paralelní algoritmy je nutno navrhovat vzhledem k paralelismu, který poskytuje HW • v případě GPU se jedná o pole SIMT m u Iti procesorů pracující nad společnou pamětí Dekompozice pro GPU • hrubozrnné rozdělení problému na části nevyžadující intenzivní komunikaci/synchronizaci • jemnozrnné rozdělení blízké vektorizaci (SIMT je ale více flexibilní) Jiří Filipovič GPU hardware a paralelismus GPU hardware Paralelismus Paměťová hierarchie Synchronizace Násobení matic 0*00000 OOOOOOOOOO OOOOOOOO OOOOOOOOOOO Hierarď lie vláken Grid Block (O, O) Block (1,0) Block (2,0) Block (1,1) Thread (0, 0) Thread (1, 0) 1 Thread (2, 0) Thread (3, 0) Thread (0,1) Thread (1,1) Thread (2,1) Thread (3,1) Thread (0, 2) Thread (1, 2) i Thread (2, 2) Thread (3, 2) Jiří Filipovič GPU hardware a paralelismus GPU hardware Paralelismus Paměťová hierarchie Synchronizace Násobení matic OOO OOÄOOOO OOOOOOOOOO OOOOOOOO OOOOOOOOOOO SIMT M uItiprocesor (G80) má jen jednu jednotku spouštějící instrukce • všech 8 SP musí provádět stejnou instrukci • nová instrukce je spuštěna každé 4 cykly • 32 thredů (tzv. warp) musí provádět stejnou instrukci Jiří Filipovič GPU hardware a paralelismus GPU hardware Paralelismus Paměťová hierarchie Synchronizace Násobení matic OOO OOÄOOOO OOOOOOOOOO OOOOOOOO OOOOOOOOOOO SIMT M uItiprocesor (G80) má jen jednu jednotku spouštějící instrukce • všech 8 SP musí provádět stejnou instrukci • nová instrukce je spuštěna každé 4 cykly • 32 thredů (tzv. warp) musí provádět stejnou instrukci A co větvení kódu? • pokud část threadů ve warpu provádí jinou instrukci, běh se serializuje • to snižuje výkon, snažíme se divergenci v rámci warpu předejít Jiří Filipovič GPU hardware a paralelismus GPU hardware Paralelismus Paměťová hierarchie Synchronizace Násobení matic OOO OOÄOOOO OOOOOOOOOO OOOOOOOO OOOOOOOOOOO SIMT M uItiprocesor (G80) má jen jednu jednotku spouštějící instrukce • všech 8 SP musí provádět stejnou instrukci • nová instrukce je spuštěna každé 4 cykly • 32 thredů (tzv. warp) musí provádět stejnou instrukci A co větvení kódu? • pokud část threadů ve warpu provádí jinou instrukci, běh se serializuje • to snižuje výkon, snažíme se divergenci v rámci warpu předejít M u Iti procesor je tedy MIMD (Multiple-Instruction Multiple-Thread) z hlediska korektnosti (téměř) a SIMT (Single-Instruction Multiple-Thread) z výkonového. Jiří Filipovič GPU hardware a paralelismus Paralelismus Paměťová hierarchie Synchronizace Násobení matic 000*000 oooooooooo oooooooo ooooooooooo SIMT rekonvergence Kompilátor označí pro divergentní kód bod rekonvergence • všechny thready mohou pokračovat až v momentě, kdy jej dosáhnou • zaručuje synchronnost threadů (jinak by pojedná divergenci nemuseli nikdy skonvergovat) • z hlediska korektnosti kódu je nutné brát tuto vlastnost v úvahu Jiří Filipovič GPU hardware a paralelismus Paralelismus Paměťová hierarchie Synchronizace Násobení matic 000*000 oooooooooo oooooooo ooooooooooo SIMT rekonvergence Kompilátor označí pro divergentní kód bod rekonvergence • všechny thready mohou pokračovat až v momentě, kdy jej dosáhnou • zaručuje synchronnost threadů (jinak by pojedná divergenci nemuseli nikdy skonvergovat) • z hlediska korektnosti kódu je nutné brát tuto vlastnost v úvahu __shared__ int s = 0; while (s != tid); s++; V díky bodu rekonvergence nad inkrementací s deadlock, pro nezávislé thready by deadlock nenastal. Jiří Filipovič GPU hardware a paralelismus GPU hardware Paralelismus Paměťová hierarchie Synchronizace Násobení matic OOO OOOOÍOO oooooooooo oooooooo ooooooooooo GPU architektura Jiří Filipovič GPU hardware a paralelismus Paralelismus Paměťová hierarchie Synchronizace Násobení matic ooooo»o oooooooooo oooooooo ooooooooooo Vlastnosti threadů Oproti CPU threadům jsou GPU thready velmi lehké (lightweight). • jejich běh může být velmi krátký (desítky instrukcí) • může (mělo by) jich být velmi mnoho • nemohou využívat velké množství prostředků Thready jsou seskupeny v blocích • ty jsou spouštěny na jednotlivých multiprocesorech • dostatečný počet bloků je důležitý pro škálovatelnost Počet threadů a thread bloků na multiprocesor je omezen. Jiří Filipovič GPU hardware a paralelismus GPU hardware Paralelismus Paměťová hierarchie OOOOOO* OOOOOOOOOO Synchronizace OOOOOOOO Násobení matic OOOOOOOOOOO Maskovaní latence pamětí Paměti mají latence • globální paměť má vysokou latenci (stovky cyklů) • registry a sdílená paměť mají read-after-write latenci Jiří Filipovič GPU hardware a paralelismus GPU hardware Paralelismus Paměťová hierarchie OOOOOO* OOOOOOOOOO Synchronizace OOOOOOOO Násobení matic OOOOOOOOOOO Maskovaní latence pamětí Paměti mají latence • globální paměť má vysokou latenci (stovky cyklů) • registry a sdílená paměť mají read-after-write latenci Maskování latence paměti je odlišné, než u CPU • žádné provádění instrukcí mimo pořadí • u většiny pamětí žádná cache Jiří Filipovič GPU hardware a paralelismus Paralelismus Paměťová hierarchie Synchronizace Násobení matic 000000» oooooooooo oooooooo ooooooooooo Maskování latence pamětí Paměti mají latence • globální paměť má vysokou latenci (stovky cyklů) • registry a sdílená paměť mají read-after-write latenci Maskování latence paměti je odlišné, než u CPU • žádné provádění instrukcí mimo pořadí • u většiny pamětí žádná cache Když nějaký warp čeká na data z paměti, je možné spustit jiný • umožní maskovat latenci paměti • vyžaduje spuštění řádově více vláken, než má GPU jader • plánování spuštění a přepínání threadů je realizováno přímo v HW bez overheadu Jiří Filipovič GPU hardware a paralelismus Paralelismus Paměťová hierarchie Synchronizace Násobení matic 000000» oooooooooo oooooooo ooooooooooo Maskování latence pamětí Paměti mají latence • globální paměť má vysokou latenci (stovky cyklů) • registry a sdílená paměť mají read-after-write latenci Maskování latence paměti je odlišné, než u CPU • žádné provádění instrukcí mimo pořadí • u většiny pamětí žádná cache Když nějaký warp čeká na data z paměti, je možné spustit jiný • umožní maskovat latenci paměti • vyžaduje spuštění řádově více vláken, než má GPU jader • plánování spuštění a přepínání threadů je realizováno přímo v HW bez overheadu Obdobná situace je v případě synchronizace. Jiří Filipovič GPU hardware a paralelismus Paralelismus Paměťová hierarchie Synchronizace Násobení matic OOOOOOO «000000000 oooooooo ooooooooooo Paměti lokální v rámci threadu Registry • nejrychlejší paměť, přímo využitelná v instrukcích • lokální proměnné v kernelu i proměnné nutné pro mezivýsledky jsou automaticky v registrech • pokud je dostatek registrů • pokud dokáže kompilátor určit statickou indexaci polí • mají životnost threadu (warpu) Jiří Filipovič GPU hardware a paralelismus Paralelismus Paměťová hierarchie Synchronizace Násobení matic OOOOOOO «000000000 oooooooo ooooooooooo Paměti lokální v rámci threadu Registry • nejrychlejší paměť, přímo využitelná v instrukcích • lokální proměnné v kernelu i proměnné nutné pro mezivýsledky jsou automaticky v registrech • pokud je dostatek registrů • pokud dokáže kompilátor určit statickou indexaci polí • mají životnost threadu (warpu) Lokální paměť • co se nevejde do registrů, jde do lokální paměti • ta je fyzicky uložena v DRAM, je tudíš pomalá a má dlouhou latenci • má životnost threadu (warpu) Jiří Filipovič GPU hardware a paralelismus GPU hardware Paralelismus Paměťová hierarchie Synchronizace Násobení matic ooooooo 0*00000000 OOOOOOOO OOOOOOOOOOO Paměť lokální v rámci bloku Sdílená paměť • u c.c. 1.x rychlá jako registry • nedojde-li ke konfliktům paměťových bank • instrukce umí využít jen jeden operand ve sdílené paměti (jinak je třeba explicitní load/store) • v C for CUDA deklarujeme pomocí __shared__ • proměnná ve sdílené paměti může mít dynamickou velikost (určenou při startu), pokud je deklarována jako extern bez udání velikosti pole • má životnost bloku Jiří Filipovič GPU hardware a paralelismus Paralelismus Paměťová hierarchie Synchronizace Násobení matic OOOOOOO oo«ooooooo OOOOOOOO OOOOOOOOOOO Sdílená paměť Deklarace statické sdílené paměti __shared__ float myArray[128] ; Dynamická alokace extern __shared__ char myArray[]; float *arrayl = (float*)myArray; int *array2 — ( intarray 1 [ 1 28] ; short *array3 = ( short array2 [256]; Vytvoří pole arrayl typu float velikosti 128, pole array2 typu int velikosti 256 a pole array3 plovoucí velikosti. Celkovou velikost je nutné specifikovat při spouštění kernelu. myKernel<«gr id , block, n>>>(); Jiří Filipovič GPU hardware a paralelismus GPU hardware Paralelismus Paměťová hierarchie Synchronizace Násobení matic ooo ooooooo ooo«oooooo oooooooo ooooooooooo Paměť lokální pro GPU Globální paměť • řádově nižší přenosová rychlost než u sdílené paměti • latence ve stovkách GPU cyklů • pro dosažení optimálního výkonu je třeba paměť adresovat zarovnaně • má životnost aplikace • u Fermi LI cache (128 byte na řádek) a L2 cache (32 byte na řádek) Lze dynamicky alokovat pomocí cudaMalloc, či staticky pomocí deklarace __c/ew'ce__ Jiří Filipovič GPU hardware a paralelismus GPU hardware Paralelismus Paměťová hierarchie Synchronizace Násobení matic ooo ooooooo oooo»ooooo oooooooo ooooooooooo Paměť lokální pro GPU Paměť konstant • pouze pro čtení • cacheována • cache-hit poskytuje rychlost jako registry (za dodržení určitých podmínek), cache-miss rychlý jako globální paměť • omezená velikost (64 KB u v současné dobe dostupných GPU) • má životnost aplikace Jiří Filipovič GPU hardware a paralelismus Paralelismus Paměťová hierarchie Synchronizace Násobení matic ooooooo ooooo«oooo oooooooo ooooooooooo Paměť konstant V deklaraci používáme __constant—, ke kopírování dat do paměti konstant slouží funkce cudaError_t cndaMemcpyToSymbol(const char ^symbol, const void *src, size_t count, size_t offset, enum cudaMemcpyKind kind) Kopíruje data ze systémové (cudaMemcpyHostToDevice) nebo globální paměti (cudaMemcpyDeviceToDevice) z src do symbol. Kopírovaný blok má velikost count bytů, do paměti symbol kopírujeme s posuvem offset. Jiří Filipovič GPU hardware a paralelismus Paralelismus ooooooo Pamětová hierarchie oooooo»ooo Synchronizace oooooooo Paměť lokální pro GPU Paměť textur • cacheovaná, 2D prostorová lokalita • pouze pro čtení (omezení kvůli cache-koherenci) • dlouhá latence • více adresovacích módů • možnost normalizace dimenzí do [0,1] • při adresaci mimo rozsah ořezávání či přetékání koordinát • možnost filtrace dat • lineární interpolace nebo nejbližší hodnota • funkcionalita je ,,zdarma" (implementováno v HW) Více detailů viz CUDA Programming Guide. i -00.0 GPU hardware a paralelismus Paralelismus Paměťová hierarchie Synchronizace Násobení matic ooooooo ooooooo«oo oooooooo ooooooooooo Paměť lokální pro systém Systémová paměť RAM • s GPU spojena přes PCI-E » virtuální adresace komplikuje přenosy mezi CPU (host) a GPU (device) pamětí • je možno alokovat tzv. page-locked oblast paměti • může redukovat celkový výkon systému • omezená velikost • data jsou po PCI-E přenášena rychleji (pro FSB) • umožňuje paralelní běh kernelu a kopírování dat • umožňuje namapovat adresovací prostor host paměti na device • umožňuje write-combining přístup (data nejsou cacheována CPU) Jiří Filipovič GPU hardware a paralelismus Paralelismus Paměťová hierarchie Synchronizace Násobení matic OOOOOOO OOOOOOOO0O OOOOOOOO OOOOOOOOOOO Page-locked paměť Namísto mallocQ použijeme pro alokaci cudaMallocHost(), pro uvolnění cudaFreeHost() • flag cudaHostAllocPortable zajistí vlastnosti page-locked paměti pro všechny (CPU) thready • flag cudaHostAllocWriteCombined vypne cacheování alokované paměti CPU • flag cudaHostAllocMapped nastaví mapování host paměti v device paměťovém prostoru Jiří Filipovič GPU hardware a paralelismus Paralelismus Paměťová hierarchie Synchronizace Násobení matic OOOOOOO OOOOOOOOO* OOOOOOOO OOOOOOOOOOO Page-locked paměť Mapovaná paměť • totéž místo má rozdílnou adresu pro device a host kód • adresu pro device získáme pomocí cudaHostGetDevicePointer() • před voláním ostatních CUDA API funkcí je zapotřebí zavolat cudaSetDeviceFlagsQ s flagem cudaDeviceMapHost Asynchronní přenosy • API funkce mají suffix Async • může se překrývat přenos dat a výpočet na CPU i přenos dat a výpočet na GPU (podrobněji až probereme streamy) Necacheovaná paměť • pomalé čtení z host kódu • rychlejší přístup z device paměti • „nevylévá" CPU cache Jiří Filipovič GPU hardware a paralelismus Paralelismus Paměťová hierarchie Synchronizace Násobení matic OOOOOOO OOOOOOOOOO »0000000 ooooooooooo Synchronizace v rámci bloku • nativní bariérová synchronizace • musí do ní vstoupit všechny thready (pozor na podmínky!) • pouze jedna instrukce, velmi rychlá, pokud neredukuje paralelismus • v C for CUDA volání __syncthreads() • Fermi rozšíření: count, and, or • komunikace přes sdílenou paměť • thready si přes ni mohou vyměňovat data • synchronizace atomickými operacemi, nebo bariérou Jiří Filipovič GPU hardware a paralelismus GPU hardware Paralelismus Paměťová hierarchie Synchronizace Násobení matic OOOOOOO OOOOOOOOOO o«oooooo OOOOOOOOOOO Atomické operace • provádí read-modify-write operace nad sdílenou nebo globální pamětí • žádná interference s ostatními thready • pro celá 32-bitová či 64-bitová (pro compute capability > 1.2) čísla (float add u c.c. > 2.0) • nad globální pamětí u zařízení s compute capability > 1.1, nad sdílenou c.c. > 1.2 • aritmetické (Add, Sub, Exch, Min, Max, lne, Dec, CAS) a bitové (And, Or, Xor) operace Jiří Filipovič GPU hardware a paralelismus Paralelismus Paměťová hierarchie Synchronizace Násobení matic ooooooo oooooooooo oo»ooooo ooooooooooo Hlasování warpu Všechny thready v jednom warpu vyhodnocují podmínku a provedou její srovnání. Dostupné u zařízení s c.c. > 1.2. int __all(int predicate); Nabývá nenulové hodnoty tehdy a jen tehdy když je nenulový predikát pro všechny thready ve warpu. int __any(int predicate); Nebývá nenulové hodnoty tehdy a jen tehdy když alespoň jeden thread ve warpu vyhodnotí predikát jako nenulový. unsigned int __ballot(int predicate); Obsahuje bitovou masku hlasování jednotlivých threadů, c.c. > 2.0 Jiří Filipovič GPU hardware a paralelismus GPU hardware Paralelismus Paměťová hierarchie OOO OOOOOOO OOOOOOOOOO Synchronizace ooo«oooo Násobení matic OOOOOOOOOOO Synchronizace paměťových operací Sdílenou paměť obvykle využíváme ke komunikaci mezi thready a nebo jako cache pro data užívaná více thready. • thready využívají data uložené jinými thready • je třeba zajistit, abychom nečetli data, která ještě nejsou k dispozici • chceme-li počkat, až jsou data k dispozici, používáme __syncthreads() Jiří Filipovič GPU hardware a paralelismus Paralelismus Paměťová hierarchie Synchronizace Násobení matic OOOOOOO OOOOOOOOOO OOOO0OOO OOOOOOOOOOO Synchronizace paměťových operací Kompilátor může optimalizovat operace se sdílenou/globální pamětí (mezivýsledky mohou zůstat v registrech) a může měnit jejich pořadí, • chceme-li se ujistit, že jsou námi ukládaná data viditelná pro ostatní, používáme __threadfence(), popř. __threadfence_block() • deklarujeme-li proměnnou jako volatile, jsou veškeré přístupy k ní realizovány přes load/store do sdílené či globální paměti • velmi důležité pokud předpokládáme implicitní synchronizaci warpu Jiří Filipovič GPU hardware a paralelismus Paralelismus Paměťová hierarchie Synchronizace Násobení matic ooooooo oooooooooo ooooo«oo ooooooooooo Synchronizace bloků Mezi bloky • globální paměť viditelná pro všechny bloky • slabá nativní podpora synchronizace • žádná globální bariéra • u novějších GPU atomické operace nad globální pamětí • globální bariéru lze implementovat voláním kernelu (jiné řešení dosti trikové) • slabé možnosti globální synchronizace znesnadňují programování, ale umožňují velmi dobrou škálovatelnost Jiří Filipovič GPU hardware a paralelismus Paralelismus Paměťová hierarchie Synchronizace Násobení matic ooooooo oooooooooo oooooo»o ooooooooooo Globální synchronizace přes atomické operace Problém součtu všech prvků vektoru • každý blok sečte prvky své části vektoru • globální bariéra • jeden blok sečte výsledky ze všech bloků Jiří Filipovič GPU hardware a paralelismus GPU hardware Paralelismus Paměťová hierarchie Synchronizace Násobení matic 000 OOOOOOO OOOOOOOOOO 0000000« OOOOOOOOOOO __device__ unsigned int count = 0; __shared__ bool isLastBlockDone; __global__ void sum(const float* array, unsigned int N, float* result) { float partialSum = calculatePartialSum(array, N); if (threadldx.x = 0) { result[blockldx.x] = partialSum; __threadfence(); unsigned int value = at omicine(&count , gridDim.x); isLastBlockDone = (value = (gridDim.x — 1)); } __syncthreads(); if (isLastBlockDone) { float totalSum = calculateTotalSum(result); if (threadldx.x = 0) { result[0] = totalSum; count = 0; } } } Jiří Filipovič GPU hardware a paralelismus Paralelismus Paměťová hierarchie Synchronizace Násobení matic OOOOOOO OOOOOOOOOO OOOOOOOO »0000000000 Násobení matic Chceme vynásobit matice A a B a výsledek uložit do C. Pro jednoduchost uvažujme matice velikosti n x n. Zápis v jazyce C for (int i — 0; i < n; i++) for (int j = 0; j < n; j++){ C[i*n + j ] = 0.0; for (int k = 0; k < n; k++) C[i*n + j] +— A[i*n + k] * B[k*n + j]; } Jiří Filipovič GPU hardware a paralelismus Paralelismus Paměťová hierarchie Synchronizace Násobení matic ooooooo oooooooooo oooooooo o«ooooooooo Paralelizace for (int i — 0; i < n; i++) for (int j = 0; j < n; j++){ C[i*n + j ] = 0.0; for (int k = 0; k < n; k++) C [ i * n + j ] += A[i*n + k] * B[k*n + j } Problém lze paralelizovat více způsoby • vybrat jeden z cyklů • vybrat dva z cyklů • paralelizovat všechny cykly Jiří Filipovič GPU hardware a paralelismus Paralelismus Paměťová hierarchie Synchronizace Násobení matic ooooooo oooooooooo oooooooo oo»oooooooo Paralelizace Paralelizace přes jeden cyklus • neškáluje dobře, nutno používat velké matice (nezapomínejme, pro dobré využití GPU potřebujeme tisíce threadů) Jiří Filipovič GPU hardware a paralelismus Paralelismus Paměťová hierarchie Synchronizace Násobení matic ooooooo oooooooooo oooooooo oo»oooooooo Paralelizace Paralelizace přes jeden cyklus • neškáluje dobře, nutno používat velké matice (nezapomínejme, pro dobré využití GPU potřebujeme tisíce threadů) Paralelizace přes dva cykly • z hlediska škálováníse zdá dobrá, počet vláken roste kvadraticky vzhledm k n Jiří Filipovič GPU hardware a paralelismus Paralelismus Paměťová hierarchie Synchronizace Násobení matic ooooooo oooooooooo oooooooo oo»oooooooo Paralelizace Paralelizace přes jeden cyklus • neškáluje dobře, nutno používat velké matice (nezapomínejme, pro dobré využití GPU potřebujeme tisíce threadů) Paralelizace přes dva cykly • z hlediska škálováníse zdá dobrá, počet vláken roste kvadraticky vzhledm k n Paralelizace přes vnitřní cyklus • nevhodná, nutná serializace zápisu do C! Jiří Filipovič GPU hardware a paralelismus Paralelismus Paměťová hierarchie Synchronizace Násobení matic ooooooo oooooooooo oooooooo oo»oooooooo Paralelizace Paralelizace přes jeden cyklus • neškáluje dobře, nutno používat velké matice (nezapomínejme, pro dobré využití GPU potřebujeme tisíce threadů) Paralelizace přes dva cykly • z hlediska škálováníse zdá dobrá, počet vláken roste kvadraticky vzhledm k n Paralelizace přes vnitřní cyklus • nevhodná, nutná serializace zápisu do C! Jako nejvhodnější se tedy jeví paralelizovat cykly jdoucí přes / a j. Jiří Filipovič GPU hardware a paralelismus Paralelismus Paměťová hierarchie Synchronizace Násobení matic ooooooo oooooooooo oooooooo ooo«ooooooo První kernel S výhodou můžeme využít možnosti uspořádání bloku a mřížky jako 2D pole. __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; } Povšimněte si nápadné podobnosti s matematickým zápisem -paralelní verze je intuitivnější, než sériová! Jiří Filipovič GPU hardware a paralelismus GPU hardware Paralelismus Paměťová hierarchie Synchronizace Násobení matic OOOOOOO OOOOOOOOOO OOOOOOOO OOOO^OOOOOO Výkon Jaký bude mít naše implementace výkon? Jiří Filipovič GPU hardware a paralelismus GPU hardware Paralelismus Paměťová hierarchie Synchronizace Násobení matic OOOOOOO OOOOOOOOOO OOOOOOOO OOOO^OOOOOO Výkon Jaký bude mít naše implementace výkon? Uvažujme kartu GeForce GTX 280 • pro problém násobení matic využitelných 622GFIops • propustnost paměti 142 GB/s Jiří Filipovič GPU hardware a paralelismus GPU hardware Paralelismus Paměťová hierarchie Synchronizace Násobení matic OOOOOOO OOOOOOOOOO OOOOOOOO OOOO^OOOOOO Výkon Jaký bude mít naše implementace výkon? Uvažujme kartu GeForce GTX 280 • pro problém násobení matic využitelných 622GFIops • propustnost paměti 142 GB/s Flop-to-word ratio naší implementace • v jednom kroku cyklu přes k načítáme 2 floaty (jedno číslo z matice A, jedno z B) a provádíme dvě aritmetické operace • jedna aritmetická operace připadne na přenos jednoho floatu (4 bytů) • globální paměť má propustnost 35.5 miliardy floatů za sekundu, pokud jeden warp přenáší jeden float z jedné matice a 16 floatů z druhé, můžeme dosáhnout výkonu cca 66.8GFIops • 66.8GFIops je velmi daleko od 622GFIops Jiří Filipovič GPU hardware a paralelismus GPU hardware Paralelismus Paměťová hierarchie Synchronizace Násobení matic OOOOOOO OOOOOOOOOO OOOOOOOO OOOOOÄOOOOO Co s tím? Narazili jsme na výkon globální paměti. GPU poskytuje rychlejší paměti, můžeme je využít? Jiří Filipovič GPU hardware a paralelismus GPU hardware Paralelismus Paměťová hierarchie Synchronizace Násobení matic OOOOOOO OOOOOOOOOO OOOOOOOO OOOOOÄOOOOO Co s tím? Narazili jsme na výkon globální paměti. GPU poskytuje rychlejší paměti, můžeme je využít? Pro výpočet jednoho prvku C musíme načíst řádek z /4 a sloupec z B, které jsou v globální paměti. Jiří Filipovič GPU hardware a paralelismus GPU hardware Paralelismus Paměťová hierarchie Synchronizace Násobení matic OOOOOOO OOOOOOOOOO OOOOOOOO OOOOOÄOOOOO Co s tím? Narazili jsme na výkon globální paměti. GPU poskytuje rychlejší paměti, můžeme je využít? Pro výpočet jednoho prvku C musíme načíst řádek z/la sloupec z B, které jsou v globální paměti. Je opravdu nutné dělat to pro každý prvek C zvlášť? • pro všechny prvky C ve stejném řádku načítáme stejný řádek A • pro všechny prvky C ve stejném sloupci načítáme stejný sloupec B • můžeme některá data načíst jednou z globální paměti do sdílené a následně je opakovaně číst z rychlejší sdílené paměti Jiří Filipovič GPU hardware a paralelismus GPU hardware Paralelismus Paměťová hierarchie Synchronizace Násobení matic OOO OOOOOOO OOOOOOOOOO OOOOOOOO OOOOOO^OOOO Přístup po blocích Budeme-li přistupovat k matici po blocích, můžeme amortizovat přenosy z globální paměti: • počítáme část matice C o velikosti a x a • iterativně načítáme bloky stejné velikosti z matic A a B do sdílené paměti • tyto bloky spolu násobíme a výsledek přičítáme do C • poměr aritmetických operací k přeneseným floatům je a-násobný Přirozené mapování na paralelismus v GPU • jednotlivé bloky threadů budou počítat bloky matice C • mají společnou sdílenou paměť • rychle se synchronizují • mezi bloky není třeba žádné synchronizace Jiří Filipovič GPU hardware a paralelismus GPU hardware Paralelismus Paměťová hierarchie Synchronizace Násobení matic OOO OOOOOOO OOOOOOOOOO OOOOOOOO ooooooo«ooo Přístup po blocích Jak velké bloky zvolit? • jsme omezeni velikostí sdílené paměti • jsme omezeni počtem threadů, který může běžet na GPU • necháme-li jeden thread počítat jeden prvek C, jeví se jako rozumná velikost bloku 16 x 16 • jedná se o násobek velikosti warpu • jeden blok bude mít únosných 256 threadů • jeden blok spotřebuje 2 KB sdílené paměti • paměť nebude zásadně omezovat výkon (16 • 35.5 — 568 GFIops, což je již poměrně blízko hodnotě 622GFIops) Jiří Filipovič GPU hardware a paralelismus Paralelismus Paměťová hierarchie Synchronizace Násobení matic OOOOOOO OOOOOOOOOO OOOOOOOO 00000000*00 Algoritmus Schéma algoritmu • každý blok threadů bude mít pole As a Bs ve sdílené paměti • iterativně se budou násobit bloky matic A a B, výsledek bude každý thread kumulovat v proměnné Csub • thready v bloku společně načtou bloky do As a Bs • každý thread vynásobí bloky v As a Bs pro jeho prvek výsledné matice v Csub • každý thread uloží jeho prvek matice do globální paměti C Pozor na synchronizaci • než začneme násobit jednotlivé bloky, musí být kompletně načteny » než začneme znovunačítat bloky, musí být dokončeno násobení s původními daty Jiří Filipovič GPU hardware a paralelismus Paralelismus Paměťová hierarchie Synchronizace Násobení matic OOOOOOO OOOOOOOOOO OOOOOOOO OOOOOOOOOÍO Druhý kernel _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_SIZE][BLOCK_SIZE]; __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE]; float Csub — O.Of; for (int b = 0; b < n/BL0CK_SIZE; b++){ As[ty][tx] — A[(ty + by*BL0CK_SIZE)* n + b*BL0CK_SIZE+tx]; Bs[ty][tx] — B[(ty + b*BL0CK_SIZE)* n + bx*BLOCK_SIZE+tx]; __syncthreads(); for (int k = 0; k < BL0CK_SIZE; k++) Csub +— As[ty][k]*Bs[k][tx]; __syncthreads(); } C[(ty + by*BL0CK)*n + bx*BL0CK_SIZE+tx] — Csub; } Jiří Filipovič GPU hardware a paralelismus Paralelismus Paměťová hierarchie Synchronizace Násobení matic OOOOOOO OOOOOOOOOO OOOOOOOO OOOOOOOOOO* Výkon • teoretické omezení první verze kernelu je 66.8GFIops, naměřený výkon 36.6GFIops • teoretické omezení druhé verze kernelu je 568GFIops, naměřený výkon 198GFIops • jak se přiblížit maximálnímu výkonu karty? • je třeba znát podrobněji HW a jeho omezení a podle toho navrhovat algoritmy • látka na další přednášku :-) Jiří Filipovič GPU hardware a paralelismus