Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooooo ooooooooo oooooo oooooooooooo oooooo oooooooo Základy obecných výpočtů na GPU Jiří Filipovič jaro 2014 Základy obecných výpočtů na GPU Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic ♦OOOOO OOOOOOO OOOOOOOOO OOOOOO OOOOOOOOOOOO OOOOOO OOOOOOOO Motivace - Moorův zákon loorův zákon Počet tranzistorů na jednom čipu se přibližně každých 18 měsíců zdvojnásobí. Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic ♦OOOOO OOOOOOO OOOOOOOOO OOOOOO OOOOOOOOOOOO OOOOOO OOOOOOOO Motivace - Moorův zákon loorův zákon Počet tranzistorů na jednom čipu se přibližně každých 18 měsíců zdvojnásobí. Adekvátní růst výkonu je zajištěn: • dříve zvyšováním frekvence, instrukčním paralelismem, out-of-order spouštěním instrukcí, vyrovnávacími pamětmi atd. • dnes vektorovými instrukcemi, zmnožováním jader Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic o«oooo ooooooo ooooooooo oooooo oooooooooooo oooooo oooooooo Motivace - grafické výpočty datově paralelní • provádíme stejné výpočty pro různé vertexy, pixely, ... • datově paralelní procesory mají vyšší koncentraci ALU, přináší tak vyšší teoretický výkon předdefinované funkce programovatelné funkce • specifické grafické efekty • GPU se stávají stále více programovatelnými • díky tomu lze zpracovávat i jiné, než grafické úlohy Základy obecných výpočtů na GPU i -00.0 •0 0.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic 0000*0 ooooooo ooooooooo oooooo oooooooooooo oooooo oooooooo Motivace - uplatnění Využití GPU pro obecné výpočty je dynamicky se rozvíjející oblast s širokou škálou aplikací Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic 0000*0 ooooooo ooooooooo oooooo oooooooooooo oooooo oooooooo Motivace - uplatnění Využití GPU pro obecné výpočty je dynamicky se rozvíjející oblast s širokou škálou aplikací • vysoce náročné vědecké výpočty • výpočetní chemie • fyzikální simulace • zpracování obrazů • a mnohé další... Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU 0000*0 ooooooo ooooooooo Motivace - uplatnění Využití GPU pro obecné výpočty je dynamicky se rozvíjející oblast s širokou škálou aplikací • vysoce náročné vědecké výpočty • výpočetní chemie • fyzikální simulace • zpracování obrazů • a mnohé další... • výpočetně náročné aplikace pro domácí uživatele • kódování a dekódování multimediálních dat • herní fyzika • úprava obrázků, 3D rendering • atd... Jiří Filipovič Základy obecných výpočtů na GPU 6/55 CUDA Demonstrační kód oooooo oooooooooooo CUDA podrobně oooooo Transpozice matic oooooooo Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic ooooo* ooooooo ooooooooo oooooo oooooooooooo oooooo oooooooo Motivace - obsah přednášky Pro plné porozumění architektuře, paralelizaci a optimalizaci pro GPU jedna přednáška nestačí • detailněji v PV197 Přednáška poskytuje základní přehled • jak vypadá architektura GPU a v čem se liší od klasických CPU • jaké druhy algoritmů běží na GPU efektivně • jak se GPU programují Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO «000000 ooooooooo oooooo oooooooooooo oooooo oooooooo Architektura GPU CPU vs. GPU • jednotky jader vs. desítky multiprocesorů • out of order vs. in order • MIMD, SIMD pro krátké vektory vs. SIMT pro dlouhé vektory • velká cache vs. malá cache, často pouze pro čtení GPU používá více tranzistorů pro výpočetní jednotky než pro cache a řízení běhu => vyšší výkon, méně univerzální Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO 0*00000 OOOOOOOOO OOOOOO OOOOOOOOOOOO OOOOOO oooooooo Architektura GPU CPU GPU Jiří Filipovič Základy obecných výpočtů na GPU 9/55 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo oo«oooo ooooooooo OOOOOO oooooooooooo OOOOOO oooooooo Architektura GPU V rámci systému: • koprocesor s dedikovanou pamětí • asynchronní běh instrukcí • připojen k systému přes PCI-E Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO 000*000 ooooooooo oooooo oooooooooooo oooooo oooooooo Procesor G80 G80 první CUDA procesor obsahuje 16 m u Iti procesorů m u Iti procesor • 8 skalárních procesorů • 2 jednotky pro speciální funkce • až 768 threadů • HW přepínání a plánování threadů • thready organizovány po 32 do warpů • SIMT • nativní synchronizace v rámci multiprocesoru Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO OOOOÍOO ooooooooo oooooo oooooooooooo oooooo oooooooo Paměťový model G80 Paměťový model • 8192 registrů sdílených mezi všemi thready multiprocesoru • 16 KB sdílené paměti • lokální v rámci multiprocesoru • stejně rychlá jako registry (za dodržení určitých podmínek) • paměť konstant • cacheovaná, pouze pro čtení • paměť pro textury • cacheovaná, 2D prostorová lokalita, pouze pro čtení • globální paměť • pro čtení i zápis, necacheovaná • přenosy mezi systémovou a grafickou pamětí přes PCI-E Jiří Filipovič Základy obecných výpočtů na GPU 12 /55 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooo*o ooooooooo OOOOOO oooooooooooo OOOOOO oooooooo Procesor G80 Device Multiprocessor N Multiprocessor 2 Multiprocessor 1 Jiří Filipovič Základy obecných výpočtů na GPU 13/55 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo oooooo* ooooooooo oooooo oooooooooooo oooooo oooooooo Další vývoj Procesory odvozené od G80 • double-precision výpočty • relaxovány pravidla pro efektivní přístup ke globální paměti • lepší možnosti synchronizace (atomické operace) Fermi • vyšší paralelizace na úrovni multiprocessoru • konfigurovatelná LI a sdílená L2 cache • plochý adresní prostor • lepší přesnost v plovoucí řádové čárce • paralelní běh kernelů Kepler • vyšší paralelizace na úrovni multiprocessoru • omezení cacheovaní • dynamický paralelismus • efektivní komunikace v rámci warpu Jiří Filipovič Základy obecných výpočtů na GPU 14/55 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO OOOOOOO «00000000 oooooo oooooooooooo oooooo oooooooo Srovnání teoretické rychlosti GPU a CPU Teoretická maxima • GPU má cca lOx rychlejší aritmetiku • GPU má cca 5x vyšší propustnost paměti • zajímavé pro mnohé problémy (budu čekat na výsledky simulace měsíc nebo rok? pojede mi video na 3 nebo 30fps?) Některé publikace ukazují lOOx i lOOOx zrychlení • v pořádku, je-li interpretováno jako zrychlení oproti produkčnímu SW (ten nemusí být perfektně optimalizovaný) • interpretováno jako srovnání CPU a GPU zpravidla nesmysl Srovnáváme-li přínos GPU oproti CPU, musíme uvažovat efektivní implementaci pro obě platformy. Jiří Filipovič Základy obecných výpočtů na GPU 15 /55 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooooo o»ooooooo oooooo oooooooooooo oooooo oooooooo Srovnání teoretické rychlosti GPU a CPU V praxi máme však často sériový CPU kód • běh v jednom vlákně znamená až 16x zpomalení (16-jádrové CPU) • absence vektorizace znamená až 4x zpomalení (32-bit operace u SSE instrukcí), 8x u AVX instrukcí Oproti sériové implementaci tedy můžeme kód paralelizací a vektorizací zrychlit • 32x pro čtyřjádrové CPU s AVX nebo osmijádrové s SSE GPU akcelerací pak • cca 300 x Vektorizace a paralelizace pro CPU je však programátorskou náročností srovnatelná s GPU akcelerací. Jiří Filipovič Základy obecných výpočtů na GPU 16/55 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooooo oo«oooooo oooooo oooooooooooo oooooo oooooooo Teoretické vs. dosažitelné zrychlení Výkonový odstup GPU může být vyšší • jednotky pro speciální funkce, operace na texturách • SIMT pružnější než SIMD • neduhy SMP (omezení škálování propustnosti paměti, „vytloukání řádků cache") Stejně jako nižší • nedostatek paralelismu • příliš vysoký overhead • nevhodný algoritmus pro GPU architekturu Dále se podíváme, jak rozlišit, jestli je nebo naopak není váš algoritmus vhodný pro GPU. Jiří Filipovič Základy obecných výpočtů na GPU 17/55 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooooo ooo«ooooo oooooo oooooooooooo oooooo oooooooo Paralelizace Sčítání vektorů • jednoduché datově-paralelní vyjádření • žádná synchronizace • potřebujeme velké vektory Game of Life • co chceme paralelizovat? Game of Life - zjištění nového stavu hry • pro větší herní plochy dostatek paralelismu • jednoduchá synchronizace Game of Life - zjištění stavu buňky po n krocích • inherentně sekvenční? (Game of Life je P-complete, P = A/C) • neznáme paralelní algoritmus Jiří Filipovič Základy obecných výpočtů na GPU 18/55 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooooo oooo«oooo oooooo oooooooooooo oooooo oooooooo Paralelizace Redukce • na první pohled může vypadat sekvenčně • ve skutečnosti realizovatelná v logn krocích • často je třeba nedržet se sekvenční verze a zamyslet se nad paralelizací problému (ne sekvenčního algoritmu) Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooooo ooooo«ooo oooooo oooooooooooo oooooo oooooooo Paralelizace Problém nalezení povodňové mapy • máme výškovou mapu terénu, přítok vody, a chceme zjistit, jaká oblast se zatopí • sekvenčnost dána rozléváním vody • je snadné najít úlohově-paralelní algoritmus, datově-paralelní už tak ne • periodická aktualizace stavu každého bodu mapy o aktualizace omezená jen na hranice vodní plochy (šetří procesory) • rozlévání vody zametači přímkou (vhodnější pro GPU, jednodušší synchronizace) • hledání souvislých oblastí a jejich spojování (odstraňuje sekvenčnost rozlévání) • vždy práce navíc oproti sekvenční/úlohově-paralelní verzi • úkol PV197 na podzim 2010, výkon odevzdaných implementací se lišil o 4 řády í!) < ! ► l -00.0 Jiří Filipovič Základy obecných výpočtů na GPU 20/55 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO OOOOOOO OOOOOOÍOO oooooo oooooooooooo oooooo oooooooo Divergence kódu Divergence kódu • serializace, divergují-li thready uvnitř warpu • nalezení nedivergujícího algoritmu může být snadné • redukce • ale také může prakticky znemožnit akceleraci některých jinak dobře paralelizovatelných algoritmů • mnoho nezávislých stavových automatů, nepravidelné datové struktury • nutnost zamyslet se nad výrazně odlišným algoritmem pro daný problém Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO OOOOOOO 0000000*0 oooooo oooooooooooo oooooo oooooooo Divergence přístupu do paměti Divergence přístupu do paměti • není-li do paměti přistupováno po souvislých blocích v rámci warpu, snižuje se její propustnost • často složitě překonatelný problém • průchod obecného grafu • může vyžadovat využití odlišných datových struktur • práce s řídkými maticemi • u rigidnějších struktur si lze často pomoci on-chip pamětí • transpozice matic Jiří Filipovič Základy obecných výpočtů na GPU 22 /55 Motivace Architektura GPU oooooo ooooooo Latence GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOOOO* OOOOOO OOOOOOOOOOOO OOOOOO OOOOOOOO GPU je dnes často propojena se zbytkem systému přes PCI-E • kopírování vstupů/výstupů je relativně pomalé • akcelerovaný algoritmus musí provádět dostatečné množství aritmetiky na přenášená data • násobení matic je vhodné (0(n3) operací na 0(n2) dat) • sčítání vhodné není (0(n2) operací na 0(n2) dat), může být však součástí většího problému Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO OOOOOOO OOOOOOOOO «00000 oooooooooooo oooooo oooooooo CUDA CUDA (Compute Unified Device Architecture) • architektura pro paralelní výpočty vyvinutá firmou NVIDIA • poskytuje nový programovací model, který umožňuje efektivní implementaci obecných výpočtů na GPU • je možné použít ji s více programovacími jazyky OpenCL Fortran C++ CUDA Architecture Jiří Filipovič Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooooo ooooooooo o«oooo oooooooooooo oooooo oooooooo C for CUDA C for CUDA přináší rozšíření jazyka C pro paralelní výpočty • explicitně oddělen host (CPU) a device (GPU) kód • hierarchie vláken • hierarchie pamětí • synchronizační mechanismy • API Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooooo ooooooooo oocooo oooooooooooo oooooo oooooooo Hierarchie vláken Hierarchie vláken • vlákna jsou organizována do bloků • bloky tvoří mřížku • problém je dekomponován na podproblémy, které mohou být prováděny nezávisle paralelně (bloky) » jednotlivé podproblémy jsou rozděleny do malých částí, které mohou být prováděny kooperativně paralelně (thready) • dobře škál uje Jiří Filipovič Základy obecných výpočtů na GPU 26/55 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooooo ooooooooo oooaoo oooooooooooo oooooo oooooooo Hierarchie vláken Grid Block (0,0) 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č Základy obecných výpočtů na GPU 27/55 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO OOOOOOO OOOOOOOOO OOOOCO OOOOOOOOOOOO OOOOOO OOOOOOOO Hierarchie pamětí Více druhů pamětí • rozdílná viditelnost • rozdílný čas života • rozdílné rychlosti a chování • přináší dobrou škálovatelnost Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO ooooooo ooooooooo ooooo* oooooooooooo oooooo oooooooo Hierarchie pamětí Jiří Filipovič Základy obecných výpočtů na GPU 29/55 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO OOOOOOO OOOOOO «00000000000 OOOOOO Příklad - součet vektorů Chceme sečíst vektory a a b a výsledek uložit do vektoru c. m -OQ.O Jiří Filipovič Základy obecných výpočtů na GPU 30/55 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO OOOOOOO OOOOOOOOO OOOOOO «00000000000 OOOOOO OOOOOOOO Příklad - součet vektorů Chceme sečíst vektory a a b a výsledek uložit do vektoru c. Je třeba najít v problému paralelismus. Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO OOOOOOO OOOOOOOOO OOOOOO «00000000000 oooooo oooooooo Příklad - součet vektorů Chceme sečíst vektory a a b a výsledek uložit do vektoru c. Je třeba najít v problému paralelismus. Sériový součet vektorů: for (int i — 0; i < N; i++) c ľ i 1 = aíil + b[il; Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO OOOOOOO OOOOOOOOO OOOOOO «00000000000 oooooo oooooooo Příklad - součet vektorů Chceme sečíst vektory a a b a výsledek uložit do vektoru c. Je třeba najít v problému paralelismus. Sériový součet vektorů: for (int i — 0; i < N; i++) c[i] = a[i] + b[i]; Jednotlivé iterace cyklu jsou na sobě nezávislé paralelizovat, škáluje s velikostí vektoru. lze je Základy obecných výpočtů na GPU Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO OOOOOOO OOOOOOOOO OOOOOO «00000000000 oooooo oooooooo Příklad - součet vektorů Chceme sečíst vektory a a b a výsledek uložit do vektoru Je třeba najít v problému paralelismus. Sériový součet vektorů: for (int i — 0; i < N; i++) c[i] = a[i] + b[i]; Jednotlivé iterace cyklu jsou na sobě nezávislé - lze je paralelizovat, škáluje s velikostí vektoru, i-tý thread sečte i-té složky vektorů: c[i] = a[i] + b[i]; Jak zjistíme, kolikátý jsme thread? Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO OOOOOOO OOOOOOOOO OOOOOO 0»0000000000 OOOOOO OOOOOOOO Hierarchie vláken Grid Block (0,0) 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č Základy obecných výpočtů na GPU 31 /55 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooooo ooooooooo oooooo oo«ooooooooo oooooo oooooooo Identifikace vlákna a bloku C for CUDA obsahuje zabudované proměnné: • threadldx.jx, y, z} udává pozici threadu v rámci bloku • blockDim. (x, y, z} udává velikost bloku • blockldx.(x, y, z} udává pozici bloku v rámci mřížky (zje vždy 1) • gridDim.jx, y, z} udává velikost mřížky (zje vždy 1) Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO OOOOOOO OOOOOOOOO OOOOOO 000*00000000 OOOOOO OOOOOOOO Příklad - součet vektorů Vypočítáme tedy globální pozici threadu (mřížka i bloky jsou jednorozměrné): Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooooo ooooooooo oooooo ooo»oooooooo oooooo oooooooo Příklad - součet vektorů Vypočítáme tedy globální pozici threadu (mřížka i bloky jsou jednorozměrné): int i = blockldx.x*blockDim.x + threadldx.x; Základy obecných výpočtů na GPU Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooooo ooooooooo oooooo ooo»oooooooo oooooo oooooooo Příklad - součet vektorů Vypočítáme tedy globální pozici threadu (mřížka i bloky jsou jednorozměrné): int blockldx.x*blockDim. threadldx.x; Celá funkce pro paralelní součet vektorů: _global__ void addvec(float *a, float *b , float * c ) { int i = blockldx.x*blockDim.x + threadldx.x; cfil = afil + bfil; Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooooo ooooooooo oooooo ooo»oooooooo oooooo oooooooo Příklad - součet vektorů Vypočítáme tedy globální pozici threadu (mřížka i bloky jsou jednorozměrné): int i = blockldx.x*blockDim.x + threadldx.x; Celá funkce pro paralelní součet vektorů: __global__ void addvec(float *a, float *b, float *c){ int i = blockldx.x*blockDim.x + threadldx.x; c[i] = a[i] + b[i]; } Funkce definuje tzv. kernel, při volání určíme, kolik threadů a v jakém uspořádání bude spuštěno. Jiří Filipovič Základy obecných výpočtů na GPU 33/55 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooooo ooooooooo oooooo oooo«ooooooo oooooo oooooooo Kvantifikátory typů funkcí Syntaxe C je rozšířena o kvantifikátory, určující, kde se bude kód provádět a odkud půjde volat: • __device__ funkce je spouštěna na device (GPU), lze volat jen z device kódu • __global__ funkce je spouštěna na device, lze volat jen z host (CPU) kódu • __host__ funkce je spouštěna na host, lze ji volat jen z host • kvantifikátory __host__ a __device__ lze kombinovat, funkce je pak kompilována pro obojí Jiří Filipovič Základy obecných výpočtů na GPU 34/55 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO OOOOOOO OOOOOOOOO OOOOOO ooooo«oooooo OOOOOO oooooooo Ke kompletnímu výpočtu je třeba: Jiří Filipovič Základy obecných výpočtů na GPU 35 /55 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO OOOOOOO OOOOOOOOO OOOOOO ooooo«oooooo OOOOOO oooooooo Ke kompletnímu výpočtu je třeba: • alokovat paměť pro vektory, naplnit je daty Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooooo ooooooooo oooooo ooooo«oooooo OOOOOO oooooooo Ke kompletnímu výpočtu je třeba: • alokovat paměť pro vektory, naplnit je daty • alokovat paměť na GPU Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO OOOOOOO OOOOOOOOO OOOOOO ooooo«oooooo OOOOOO oooooooo Ke kompletnímu výpočtu je třeba: • alokovat paměť pro vektory, naplnit je daty • alokovat paměť na GPU • zkopírovat vektory a a b na GPU Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooooo ooooooooo oooooo ooooo«oooooo oooooo oooooooo Ke kompletnímu výpočtu je třeba: • alokovat paměť pro vektory, naplnit je daty • alokovat paměť na GPU • zkopírovat vektory a a b na GPU • spočítat vektorový součet na GPU Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooooo ooooooooo oooooo ooooo«oooooo oooooo oooooooo Ke kompletnímu výpočtu je třeba: • alokovat paměť pro vektory, naplnit je daty • alokovat paměť na GPU • zkopírovat vektory a a b na GPU • spočítat vektorový součet na GPU • uložit výsledek z GPU paměti do c Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO OOOOOOO OOOOOOOOO OOOOOO OOOOO^OOOOOO oooooo oooooooo Ke kompletnímu výpočtu je třeba: • alokovat paměť pro vektory, naplnit je daty • alokovat paměť na GPU • zkopírovat vektory a a b na GPU • spočítat vektorový součet na GPU • uložit výsledek z GPU paměti do c • použít výsledek v c :-) Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU oooooo ooooooo ooooooooo Příklad - součet vektorů CPU kód naplní a a b, vypíše c: (f include Sdefine N 64 int main(){ float a[N] , b[N] , c [N] ; for (int i = 0; i < N; i++) a[i] = b[i] = i; // zde bude kód provádějící výpočet na GPU for (int i = 0; i < N; i++) printf("%f , " , c[i ] ); return 0; } Jiří Filipovič Základy obecných výpočtů na GPU 36/55 CUDA Demonstrační kód oooooo oooooo«ooooo CUDA podrobně OOOOOO Transpozice matic oooooooo Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO OOOOOOO OOOOOOOOO OOOOOO 0000000*0000 oooooo oooooooo Správa GPU paměti Paměť je třeba dynamicky alokovat. cudaMalloc(void** devPtr, size_t count); Alokuje paměť velikosti count, nastaví na ni ukazatel devPtr. Uvolnění paměti: cudaFree(void* devPtr); Kopírování paměti: cudaMemcpy(void* dst, const void* src, size_t count, enum cudaMemcpyKind kind); Kopíruje count byte z src do dst, kind určuje, o jaký směr kopírování se jedná (např. cudaMemcpyHostToDevice, nebo cudaMemcpyDevice ToHost). Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU oooooo ooooooo ooooooooo Příklad - součet vektorů Alokujeme paměť a přeneseme data: float *d_a, *d_b, *d_c; cudaMalloc((void**)&d_a, N*sizeof(*d_a)); cud aM alloc (( vo id **)&id_b , N*sizeof(*d_b)); cudaMalloc((void**)&d_c, N*sizeof(*d_c)); cudaMemcpy(d_a, a, N*sizeof(*d_a), cudaMemcpyHostToDevice); cudaMemcpy(d_b, b, N*sizeof(*d_b), CudaMemcpyHostToDevice); // zde bude spuštěn kernel cudaMemcpy(c, d_c, N*sizeof(*c), cudaMemcpyDeviceToHost); cudaFree(d_a); cudaFree(d_b); cudaFree(d_c); Jiří Filipovič Základy obecných výpočtů na GPU 38/55 CUDA Demonstrační kód OOOOOO oooooooo«ooo CUDA podrobně OOOOOO Transpozice matic OOOOOOOO Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooooo ooooooooo oooooo ooooooooo«oo oooooo oooooooo Příklad - součet vektorů Spuštění kernelu: • kernel voláme jako funkci, mezi její jméno a argumenty vkládáme do trojitých špičatých závorek velikost mřížky a bloku • potřebujeme znát velikost bloků a jejich počet • použijeme ID blok i mřížku, blok bude pevné velikosti • velikost mřížky vypočteme tak, aby byl vyřešen celý problém násobení vektorů Pro vektory velikosti dělitelné 32: Sdefine BLOCK 32 addvec«(d_a , d_b , d_c ) ; Jak řešit problém pro obecnou velikost vektoru? Jiří Filipovič Základy obecných výpočtů na GPU 39/55 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooooo ooooooooo oooooo oooooooooo»o oooooo oooooooo Příklad - součet vektorů Upravíme kód kernelu: __global__ void addvec(float *a, float *b, float *c, int n){ int i = blockldx.x*blockDim.x + threadldx.x; if (i < n) c[i] = a[i] + b[i]; } A zavoláme kernel s dostatečným počtem vláken: addvec«(d_a , d_b , d_c , N) ; Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO OOOOOOO OOOOOOOOO OOOOOO OOOOOOOOOOO* oooooo oooooooo Příklad - spuštění Nyní už zbývá jen kompilace :-). nvcc -I/usr/local/cuda/include -L/usr/local/cuda/lib -lcudart \ -o vecadd vecadd.cu Kde s CUDA pracovat? • vlastní stroj: stáhněte a nainstalujte CUDA toolkit a SDK z developer.nvidia.com • windowsí stanice v učebnách (titan) • ke vzdálené práci s high-end GPU: na přání Základy obecných výpočtů na GPU Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO OOOOOOO OOOOOOOOO OOOOOO OOOOOOOOOOOO «00000 oooooooo 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) Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO OOOOOOO OOOOOOOOO OOOOOO OOOOOOOOOOOO «00000 OOOOOOOO 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č Základy obecných výpočtů na GPU 42 /55 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooooo ooooooooo oooooo oooooooooooo oaoooo oooooooo Paměť lokální v rámci bloku Sdílená paměť • organizována do bank umožňujících paralelní přístup • u rodiny G80 rychlá jako registry • nedojde-li ke konfliktům pamětových bank • instrukce umí využít jen jeden operand ve sdílené paměti (jinak je třeba explicitní load/store) • u novějších GPU ve srovnání s registry pomalejší • v C for CUDA deklarujeme pomocí __shared__ • má životnost bloku Jiří Filipovič Základy obecných výpočtů na GPU 43/55 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO OOOOOOO OOOOOOOOO OOOOOO OOOOOOOOOOOO 00*000 oooooooo 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 sdružené • 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č Základy obecných výpočtů na GPU 44/55 Motivace Architektura GPU OOOOOO OOOOOOO Ostatní paměti Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOOOOO OOOOOO OOOOOOOOOOOO OOOÄOO OOOOOOOO • paměť konstant • texturová paměť • systémová paměť Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO OOOOOOO OOOOOOOOO OOOOOO OOOOOOOOOOOO 0000*0 oooooooo 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 • atomické operace nad sdílenou pamětí Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooooo ooooooooo oooooo oooooooooooo ooooo* oooooooo Synchronizace bloků Mezi bloky • globální paměť viditelná pro všechny bloky • slabá nativní podpora synchronizace • žádná globální bariéra uvnitř kernelu • 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 hardware • u novějších GPU atomické operace nad globální pamětí Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO OOOOOOO OOOOOOOOO OOOOOO OOOOOOOOOOOO OOOOOO «0000000 Transpozice matic Z teoretického hlediska: • triviální problém • triviální paralelizace • jsme triviálně omezení propustností paměti (neděláme žádné flops) __global__ void mtran(float *odata, float* idata, int n){ int x = blockldx.x * blockDim.x + threadldx.x; int y = blockldx.y * blockDim.y + threadldx.y; odata[x*n + y] = idata[y*n + x]; Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO OOOOOOO OOOOOOOOO OOOOOO OOOOOOOOOOOO OOOOOO o»oooooo Výkon Spustíme-li kód na GeForce GTX 280 s použitím dostatečně velké matice 4000 x 4000, bude propustnost 5.3 GB/s. Kde je problém? Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooooo ooooooooo oooooo oooooooooooo oooooo o»oooooo Výkon Spustíme-li kód na GeForce GTX 280 s použitím dostatečně velké matice 4000 x 4000, bude propustnost 5.3 GB/s. Kde je problém? Přístup do odata je prokládaný! Modifikujeme transpozici na kopírování: odata[y*n + x] = idata[y*n + x]; a získáme propustnost 112.4 GB/s. Pokud bychom přistupovali s prokládáním i k idata, bude výsledná rychlost 2.7 GB/s. Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooooo ooooooooo oooooo oooooooooooo oooooo oo«ooooo Odstranění prokládání Matici můžeme zpracovávat po blocích • načteme po řádcích blok do sdílené paměti • uložíme do globální paměti jeho transpozici taktéž po řádcích • díky tomu je jak čtení, tak zápis bez prokládání Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooooo ooooooooo oooooo oooooooooooo oooooo oo«ooooo Odstranění prokládání Matici můžeme zpracovávat po blocích • načteme po řádcích blok do sdílené paměti • uložíme do globální paměti jeho transpozici taktéž po řádcích • díky tomu je jak čtení, tak zápis bez prokládání Jak velké bloky použít? • budeme uvažovat bloky čtvercové velikosti • pro sdružené čtení musí mít řádek bloku velikost dělitelnou 16 • v úvahu připadají bloky 16 x 16, 32 x 32 a 48 x 48 (jsme omezeni velikostí sdílené paměti) • nejvhodnější velikost určíme experimentálně Jiří Filipovič Základy obecných výpočtů na GPU 50/55 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooooo ooooooooo oooooo oooooooooooo oooooo ooo»oooo Bloková transpozice __global__ void mtran_coalesced(float *odata, float *idata, int n) __shared__ float tile[TILE_DIM][TILE_DIM]; int x — blockldx.x * TILE_DIM + threadldx.x; int y — blockldx.y * TILE_DIM + threadldx.y; int index_in = x + y*n; x — blockldx.y * TILE_DIM + threadldx.x; y — blockldx.x * TILE_DIM + threadldx.y; int index_out = x + y*n ; for (int i = 0; i < TILE_DIM; i += BL0CK_R0WS) tile[threadldx.y+i][threadldx.x] = idata[index.in+i*n]; __syncthreads(); for (int i = 0; i < TILE_DIM; i += BL0CK_R0WS) odata[index_out+i*n] = tile[threadldx.x][threadldx.y+i] ; } Jiří Filipovič Základy obecných výpočtů na GPU 51 /55 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO ooooooo ooooooooo OOOOOO oooooooooooo OOOOOO oooo«ooo Výkon Nejvyšší výkon byl naměřen při použití bloků velikosti 32 x 32, velikost thread bloku 32 x 8, a to 75.1GB/s. Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO ooooooo OOOOOOOOO OOOOOO oooooooooooo OOOOOO oooo«ooo Výkon Nej vyšší výkon byl naměřen při použití bloků velikosti 32 x 32, velikost thread bloku 32 x 8, a to 75.1GB/s. • to je výrazně lepší výsledek, nicméně stále nedosahujeme rychlosti pouhého kopírování Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooooo ooooooooo oooooo oooooooooooo oooooo oooo«ooo Výkon Nej vyšší výkon byl naměřen při použití bloků velikosti 32 x 32, velikost thread bloku 32 x 8, a to 75.1GB/s. • to je výrazně lepší výsledek, nicméně stále nedosahujeme rychlosti pouhého kopírování • kernel je však složitější, obsahuje synchronizaci • je nutno ověřit, jestli jsme narazili na maximum, nebo je ještě někde problém Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooooo ooooooooo oooooo oooooooooooo oooooo oooo«ooo Výkon Nej vyšší výkon byl naměřen při použití bloků velikosti 32 x 32, velikost thread bloku 32 x 8, a to 75.1GB/s. • to je výrazně lepší výsledek, nicméně stále nedosahujeme rychlosti pouhého kopírování • kernel je však složitější, obsahuje synchronizaci • je nutno ověřit, jestli jsme narazili na maximum, nebo je ještě někde problém • pokud v rámci bloků pouze kopírujeme, dosáhneme výkonu 94.9GB/s • něco ještě není optimální Jiří Filipovič Základy obecných výpočtů na GPU 52 /55 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO OOOOOOO OOOOOOOOO OOOOOO OOOOOOOOOOOO OOOOOO OOOOO^OO Sdílená paměť Při čtení globální paměti zapisujeme do sdílené paměti po řádcích. t ile[threadldx.y+i] [threadldx. x ] = idata[index_ in+i*n]; Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO OOOOOOO OOOOOOOOO OOOOOO OOOOOOOOOOOO OOOOOO 00000*00 Sdílená paměť Při čtení globální paměti zapisujeme do sdílené paměti po řádcích. t ile[threadldx.y+i] [threadldx.x] = idata[index_in+i*n]; Při zápisu do globální paměti čteme ze sdílené po sloupcích. odata[index_out+i*n] = tile[threadldx.x][threadldx.y+i]; To je čtení s prokládáním, které je násobkem 16, celý sloupec je tedy v jedné bance, vzniká 16-cestný bank conflict. Jiří Filipovič Základy obecných výpočtů na GPU 53/55 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooooo ooooooooo oooooo oooooooooooo oooooo OOOOO^OO Sdílená paměť Při čtení globální paměti zapisujeme do sdílené paměti po řádcích. tile[threadldx.y+i][threadldx.x] = idata[index.in+i*n]; Při zápisu do globální paměti čteme ze sdílené po sloupcích. odata[index_out+i*n] = tile[threadldx.x][threadldx.y+i] ; To je čtení s prokládáním, které je násobkem 16, celý sloupec je tedy v jedné bance, vzniká 16-cestný bank conflict. Řešením je padding: __shared__ float tile[TILE_DIM][TILE_DIM + 1]; Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic oooooo ooooooo ooooooooo oooooo oooooooooooo oooooo oooooo«o Výkon Nyní dosahuje naše implementace výkon 93.4 GB/s. • obdobný výsledek, jako při pouhém kopírování • zdá se, že výrazněji lepšího výsledku již pro danou matici nedosáhneme • pozor na různou velikost vstupních dat (tzv. partition camping, není problém u Fermi) Základy obecných výpočtů na GPU i -00.0 Motivace Architektura GPU Algoritmy a GPU CUDA Demonstrační kód CUDA podrobně Transpozice matic OOOOOO OOOOOOO OOOOOOOOO OOOOOO OOOOOOOOOOOO OOOOOO 0000000« Zhodnocení výkonu Veškeré optimalizace sloužily pouze k lepšímu přizpůsobení-se vlastnostem HW • přesto jsme dosáhli 17.6x zrychlení • při formulaci algoritmu je nezbytné věnovat pozornost hardwareovým omezením • jinak bychom se nemuseli vývojem pro GPU vůbec zabývat, stačilo by napsat dobrý CPU algoritmus... Základy obecných výpočtů na GPU