Motivace Architektura GPU C for CUDA Demonstrační kód CUDA podrobně Závěr oooo oooooooooo ooooo ooooooooooo oooooooooo oo Akcelerace výpočtů na GPU J-W i- ■ i " "V in Fihpovic jaro 2019 Jiří Filipovič Akcelerace výpočtů na GPU 1/43 >ooo Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód ooooooooooo CUDA podrobně oooooooooo Motivace - aritmetický výkon GPU Theoretical GFLOP/s 5750 5500 5250 5000 4750 4500 4250 4000 Pentium 4 Bloomfield Westmere Apr-01 Sep-02 Jan-04 May-05 Oct-06 Feb-08 Jul-09 Nov-10 Apr-12 Aug-13 Dec-14 J in Fihpovic Akcelerace výpočtů na GPU Motivace Architektura GPU C for CUDA Demonstrační kód CUDA podrobně Závěr o«oo oooooooooo ooooo ooooooooooo oooooooooo oo Motivace - paměťová propustnost GPU Theoretical GB/s 360 330 300 270 240 210 180 150 120 90 60 GeForce FX 5900 GeForce 780 Ti Ť Tesla GeForce GPU Tesla GPU- Tesla K20X GeForce GTX4S0 GeForce GTX GeForcq 6800 GT 30 GeForce 8800 GTX GeForce 7800 GTX North wood 2003 2004 2005 2006 2007 2008 2009 2010 2011 2012 2013 Jin Fihpovic Akcelerace výpočtů na GPU Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód ooooooooooo CUDA podrobně oooooooooo Motivace - náročnost programování OK, GPU jsou výkonnější, ale není jejich programování výrazně naroci a je složitější, než psát sériový skalární C/C++ kód • je to ale fér srovnání? j in Fihpovic Akcelerace výpočtů na GPU Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód ooooooooooo CUDA podrobně oooooooooo Motivace - náročnost programování OK, GPU jsou výkonnější, ale není jejich programování výrazně naroci a je složitější, než psát sériový skalární C/C++ kód • je to ale fér srovnání? Moorův zákon Počet tranzistorů umístitelných na jeden čip se zdvojnásobí každých 18 měsíců J in Fihpovic Akcelerace výpočtů na GPU Motivace oo«o Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód ooooooooooo CUDA podrobně oooooooooo Závěr oo Motivace - náročnost programování OK, GPU jsou výkonnější, ale není jejich programování výrazně náročnejší r o je složitější, než psát sériový skalární C/C++ kód... • je to ale fér srovnání? Moorův zákon Počet tranzistorů umístitelných na jeden čip se zdvojnásobí každých 18 měsíců Odpovídající růst výkonu zajišťuje: • v minulosti: zvýšení frekvence, instrukčního paralelismu, provádění instrukcí mimo pořadí atp. • dnes: vektorové instrukce, zvyšování počtu jader Jiří Filipovič Akcelerace výpočtů na GPU 4/43 Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód ooooooooooo CUDA podrobně oooooooooo Motivace - změna paradigmatu Důsledky moorova zákona: • v minulosti: změny v architektuře důležité pro vývojáře překladačů, vývojářů aplikací se příliš netýkaly • dnes: k vyuřžití plného výkonu procesorů je nutné kód paralelizovat a vektorizovat • stále úkol pro programátory, ne pro překladač • psát efektivní kód pro GPU je obdobně náročné, jako pro CPU J in Fihpovic Akcelerace výpočtů na GPU Architektura GPU •ooooooooo C for CUDA ooooo Demonstrační kód ooooooooooo CUDA podrobně oooooooooo Co dělá GPU výkonnými? Typy paralelismu • Úlohový paralelismus • dekompozice problému na úlohy, které mohou být prováděny paralelně 9 obvykle komplexní úlohy mohou provádět rozdílné činnosti • komplexnější synchronizace • vhodné pro menší počet výkonných procesorů (jader) o Datový paralelismus • paralelismus na úrovni datových struktur • obvykle stejná operace na více elementech datové struktury • umožňuje konstrukci jednodušších (a menších) procesorů j in Fihpovic Akcelerace výpočtů na GPU Architektura GPU o«oooooooo C for CUDA ooooo Demonstrační kód ooooooooooo CUDA podrobně oooooooooo Co dělá GPU výkonnými? Z pohledu programátora o některé problémy jsou spíše úlohově paralelní, některé spíše datově (průchod grafu vs. sčítání vektorů) Z pohledu hardware • procesory zpracovávající datově-paralelní úlohy mohou být jednodušší • můžeme dosáhnout vyššího aritmetického výkonu se stejnou velikostí procesoru (tj. na stejný počet tranzistorů) • jednoduché vzory přístupu do paměti umožňují konstrukci paměti s vysokou propustností j in Fihpovic Akcelerace výpočtů na GPU Motivace Architektura GPU C for CUDA Demonstrační kód CUDA podrobně Závěr OOOO OO0OOOOOOO ooooo ooooooooooo oooooooooo oo CPU GPU Jiří Filipovič Akcelerace výpočtů na GPU 8/43 Motivace oooo Architektura GPU ooo«oooooo C for CUDA ooooo Demonstrační kód ooooooooooo CUDA podrobně oooooooooo Architektura GPU CPU vs. GPU • stovky ALU v desítkách jader vs. tisíce ALU v desítkách multi procesorů • 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í Jin Fihpovic Akcelerace výpočtů na GPU Motivace oooo Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód ooooooooooo CUDA podrobně oooooooooo Architektura GPU High-end GPU: • koprocesor s dedikovanou pamětí • asynchronní běh instrukcí připojen k systému přes PCI-E Jin Fihpovic Akcelerace výpočtů na GPU Motivace oooo CUDA Architektura GPU OOOOO0OOOO C for CUDA ooooo Demonstrační kód ooooooooooo CUDA podrobně oooooooooo Závěr oo 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 9 je možné použít ji s více programovacími jazyky Jiří Filipovič Akcelerace výpočtů na GPU 11 /43 Architektura GPU OOOOOO0OOO C for CUDA ooooo Demonstrační kód ooooooooooo CUDA podrobně oooooooooo Procesor G80 G80 9 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 vláken • HW prepínania plánování vláken • vlákna organizována po 32 do warpů • SIMT • nativní synchronizace v rámci m u Iti procesoru J in Fihpovic Akcelerace výpočtů na GPU Motivace oooo Architektura GPU OOOOOOO0OO C for CUDA ooooo Paměťový model G80 Demonstrační kód ooooooooooo CUDA podrobně oooooooooo Závěr oo Paměťový model o 8192 registrů sdílených mezi všemi vlákny multiprocesoru • 16 KB sdílené paměti • lokální v rámci multi procesoru • rychlost blízká registrům (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č Akcelerace výpočtů na GPU 13/43 Motivace oooo Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód ooooooooooo CUDA podrobně oooooooooo Závěr oo Procesor G80 Device Multiprocessor N Multiprocessor 2 Multiprocessor 1 Architektura GPU 000000000« C for CUDA ooooo Demonstrační kód ooooooooooo CUDA podrobně oooooooooo Procesory CUDA architektury o double-precision výpočty • benevolentnější pravidla pro efektivní přístup ke globální paměti • LI, L2/data cache 9 navýšeny on-chip zdroje (více registrů, více vláken na MP) o lepší možnosti synchronizace • dynamický paralelismus • přístup na vstupně-výstupní porty j in Fihpovic Akcelerace výpočtů na GPU Architektura GPU oooooooooo C for CUDA •oooo Demonstrační kód ooooooooooo CUDA podrobně oooooooooo 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 9 hierarchie pamětí • synchronizační mechanismy 9 API J in Fihpovic Akcelerace výpočtů na GPU Architektura GPU oooooooooo C for CUDA o«ooo Demonstrační kód ooooooooooo CUDA podrobně oooooooooo Hierarchie vláken Hierarchie vláken o 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) o jednotlivé podproblémy jsou rozděleny do malých částí, které mohou být prováděny kooperativně paralelně (vlákna) o dobře škál uje J in Fihpovic Akcelerace výpočtů na GPU Motivace oooo Architektura GPU oooooooooo C for CUDA OO0OO Demonstrační kód ooooooooooo CUDA podrobně oooooooooo Závěr oo Hierarchie vláken Grid Block (o, o) Block (1, o) Block (2, o) Block (1,1) Thread (0, 0) Thread (1, 0) Thread (2, 0) Thread (3, 0) Thread (0,1) Thread (1,1) Thread (2,1) Thread (3,1) Thread (0, 2) Thread (1, 2) Thread (2, 2) Thread (3, 2) Jiří Filipovič Akcelerace výpočtů na GPU 18/43 Architektura GPU oooooooooo C for CUDA oooto Demonstrační kód ooooooooooo CUDA podrobně oooooooooo Hierarchie pamětí Více druhů pamětí o rozdílná viditelnost • rozdílný čas života • rozdílné rychlosti a chování • přináší dobrou škálovatelnost Jin Fihpovic Akcelerace výpočtů na GPU Motivace oooo Architektura GPU oooooooooo C for CUDA oooo* Demonstrační kód ooooooooooo CUDA podrobně oooooooooo Závěr oo Hierarchie pamětí Thread Per-thread local memory Thread Block Per-block shared memory Block (0, 2) Block (1, 2) Jiří Filipovič Akcelerace výpočtů na GPU 20/43 Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód •oooooooooo CUDA podrobně oooooooooo Příklad - součet vektorů Chceme sečíst vektory a a b a výsledek uložit do vektoru c. 4 ^ >■ < ► 4 S J in Fihpovic Akcelerace výpočtů na GPU Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód •oooooooooo CUDA podrobně oooooooooo 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. 4 ^ >■ < ► 4 -š Jin Fihpovic Akcelerace výpočtů na GPU Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód •oooooooooo CUDA podrobně oooooooooo 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]; j in Fihpovic Akcelerace výpočtů na GPU Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód •oooooooooo CUDA podrobně oooooooooo 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é - lze je paralelizovat, škáluje s velikostí vektoru. 4 ^ >■ < ► 4 -š J in Fihpovic Akcelerace výpočtů na GPU Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód •oooooooooo CUDA podrobně oooooooooo 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é - lze je paralelizovat, škáluje s velikostí vektoru, i-té vlákno sečte i-té složky vektorů: c[i] = a[i] + b[i]; Jak zjistíme, kolikáté jsme vlákno? Jin Fihpovic Akcelerace výpočtů na GPU Motivace oooo Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód O0OOOOOOOOO CUDA podrobně oooooooooo Závěr oo Hierarchie vláken Grid Block (o, o) Block (1, o) Block (2, o) Block (1,1) Thread (0, 0) Thread (1, 0) Thread (2, 0) Thread (3, 0) Thread (0,1) Thread (1,1) Thread (2,1) Thread (3,1) Thread (0, 2) Thread (1, 2) Thread (2, 2) Thread (3, 2) Jiří Filipovič Akcelerace výpočtů na GPU 22 /43 Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód oo«oooooooo CUDA podrobně oooooooooo Identifikace vlákna a bloku C for CUDA obsahuje zabudované proměnné: • threadldx.jx, y, z} udává pozici vlákna v rámci bloku • blockDim.{x, y, z} udává velikost bloku • blockldx.jx, y, z} udává pozici bloku v rámci mřížky (zje vždy 1) 9 gridDim.{x, y, z} udává velikost mřížky (z je vždy 1) j in Fihpovic Akcelerace výpočtů na GPU Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód ooo«ooooooo CUDA podrobně oooooooooo Příklad - součet vektorů Vypočítáme tedy globální pozici vlákna (mřížka i bloky jsou jednorozměrné): Jin Fihpovic Akcelerace výpočtů na GPU Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód ooo«ooooooo CUDA podrobně oooooooooo Příklad - součet vektorů Vypočítáme tedy globální pozici vlákna (mřížka i bloky jsou jednorozměrné): int i = blockldx.x*blockDim.x + threadldx.x; Jin Fihpovic Akcelerace výpočtů na GPU Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód ooo«ooooooo CUDA podrobně oooooooooo Příklad - součet vektorů Vypočítáme tedy globální pozici vlákna (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 = biockldx.x*biockDim.x + threadldx.x; c[i] = a[i] + b[i]; } Jin Fihpovic Akcelerace výpočtů na GPU Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód ooo«ooooooo CUDA podrobně oooooooooo Příklad - součet vektorů Vypočítáme tedy globální pozici vlákna (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 = biockldx.x*biockDim.x + threadldx.x; c[i] = a[i] + b[i]; } Funkce definuje tzv. kernel, při volání určíme, kolik vláken a v jakém uspořádání bude spuštěno. j in Fihpovic Akcelerace výpočtů na GPU Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód OOOO^OOOOOO CUDA podrobně oooooooooo 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í J in Fihpovic Akcelerace výpočtů na GPU Motivace Architektura GPU C for CUDA Demonstrační kód CUDA podrobně Závěr oooo oooooooooo ooooo ooooo«ooooo oooooooooo oo Ke kompletnímu výpočtu je třeba: Jiří Filipovič Akcelerace výpočtů na GPU 26/43 Motivace Architektura GPU C for CUDA Demonstrační kód CUDA podrobně Závěr oooo oooooooooo ooooo oooooooooo oooooooooo oo Ke kompletnímu výpočtu je třeba: • alokovat paměť pro vektory, naplnit je daty Jiří Filipovič rS1 Akcelerace výpočtů na GPU Ke kompletnímu výpočtu je třeba: • alokovat paměť pro vektory, naplnit je daty • alokovat paměť na GPU Jin Fihpovic rS1 Akcelerace výpočtů na GPU 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 Jiří Filipovič □ i3" Akcelerace výpočtů na GPU Motivace Architektura GPU C for CUDA Demonstrační kód CUDA podrobně Závěr oooo oooooooooo ooooo ooooo«ooooo oooooooooo oo Ke kompletnímu výpočtu je třeba: • alokovat paměť pro vektory, naplnit je daty • alokovat pamět na GPU 9 zkopírovat vektory a a b na GPU 9 spočítat vektorový součet na GPU Jiří Filipovič Akcelerace výpočtů na GPU 5 >qc^(v Ke kompletnímu výpočtu je třeba: • alokovat paměť pro vektory, naplnit je daty • alokovat paměť na GPU 9 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 4 □ ► 4 ^ >■ < ► 4 -š Jiri Fihpovic Akcelerace výpočtů na GPU Motivace Architektura GPU C for CUDA Demonstrační kód CUDA podrobně Závěr oooo oooooooooo ooooo ooooo«ooooo oooooooooo oo Ke kompletnímu výpočtu je třeba: • alokovat paměť pro vektory, naplnit je daty • alokovat paměť na GPU 9 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 :-) Při použití managed memory (od compute capability 3.0, CUDA 6.0) není třeba explicitně provádět kroky psané kurzívou. Jiří Filipovič Akcelerace výpočtů na GPU 26/43 Motivace oooo Architektura GPU oooooooooo C for CUDA ooooo Příklad - součet vektorů Demonstrační kód OOOOOO0OOOO CUDA podrobně oooooooooo Závěr oo CPU kód naplní a a b, vypíše c: #include #define N 64 int main(){ float *a, *b, *c ; cudaMallocManaged(&a, N*sizeof(*a)) cudaMallocManaged(&b, N*sizeof(*b)) cudaMallocManaged((^c , N*sizeof (* c ) ) for (int i = 0; i < N; i++) a[i] = i; b[i] = i*2; // zde bude kód provádějící výpočet na GPU for (int i = 0; i < N; i++) printf("%f, " , c[i]); cudaFree(a); cudaFree(b); cudaFree(c); return 0; } < □ ► ^ s ► < -e ► < e j in Fihpovic Akcelerace výpočtů na GPU oooo Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód OOOOOOO0OOO CUDA podrobně oooooooooo Správa GPU paměti Použili jsme managed paměť, CUDA se automaticky stará o přesuny mezi CPU a GPU. • koherence je automaticky zajištěna • k paměti nelze přistupovat, pokud běží CUDA kernel (i když ji nepoužívá) Lze použít také explicitní alokaci: cudaMalloc(void** devPtr, size_t count); cudaFree(void* devPtr); cudaMemcpy(void* dst , const void* src , size_t count, enum cudaMemcpyKind kind ) ; J in Fihpovic Akcelerace výpočtů na GPU Motivace oooo Architektura GPU oooooooooo C for CUDA ooooo Příklad - součet vektorů Demonstrační kód OOOOOOOO0OO CUDA podrobně oooooooooo Závěr oo 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: #define BLOCK 32 addvec«(a, b, c); cudaDeviceSynchronize(); Synchronizace za voláním kernelu zajistí, že výpis hodnoty v c bude proveden až po dokončení kernelu. Jiří Filipovič Akcelerace výpočtů na GPU 29/43 Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód OOOOOOOOO0O CUDA podrobně oooooooooo Příklad - součet vektorů Jak řešit problém pro obecnou velikost vektoru? Upravíme kód kernelu: __global__ void addvec(float *a, float *b , float *c , int i = biockldx.x*biockDim.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«»(a, b, c, N) ; int n){ j in Fihpovic Akcelerace výpočtů na GPU Motivace Architektura GPU C for CUDA Demonstrační kód CUDA podrobně Závěr oooo oooooooooo ooooo oooooooooo* oooooooooo oo Nyní už zbývá jen kompilace :-). nvcc -o vecadd vecadd.cu Jiří Filipovič Akcelerace výpočtů na GPU 31 /43 Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód ooooooooooo CUDA podrobně •ooooooooo Paměti lokální v rámci vlákna 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 vlákna (warpu) J in Fihpovic Akcelerace výpočtů na GPU Motivace oooo Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód ooooooooooo CUDA podrobně •ooooooooo Závěr oo Paměti lokální v rámci vlákna 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 vlákna (warpu) Lokální paměť • co se nevejde do registrů, jde do lokální paměti 9 ta je fyzicky uložena v DRAM, je tudíž pomalá a má dlouhou latenci (může být však cacheována) o má životnost vlákna (warpu) Jiří Filipovič Akcelerace výpočtů na GPU 32 /43 Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód ooooooooooo CUDA podrobně o«oooooooo Paměť lokální v rámci bloku Sdílená paměť • rychlost se blíží registrům • nedojde-li ke konfliktům paměťových bank • může vyřadovat load/store instrukce navíc o v C for CUDA deklarujeme pomocí __s/7arec/__ • 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 o má životnost bloku j in Fihpovic Akcelerace výpočtů na GPU Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód ooooooooooo CUDA podrobně OO0OOOOOOO 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 n řádek), Kepler L2, c.c. 3.5 data cache, Maxwell data cache Lze dynamicky alokovat pomocí cudaMalloc, či staticky pomocí deklarace —device— Jin Fihpovic Akcelerace výpočtů na GPU oooo Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód ooooooooooo CUDA podrobně oooooooooo Ostatní paměti • paměť konstant • texturová paměť • systémová paměť J in Fihpovic Akcelerace výpočtů na GPU oooo Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód ooooooooooo CUDA podrobně oooooooooo Synchronizace v rámci bloku nativní bariérová synchronizace • musí do ní vstoupit všechna vlákna (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 j in Fihpovic Akcelerace výpočtů na GPU Motivace oooo Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód ooooooooooo CUDA podrobně OOOOO0OOOO Atomické operace 9 provádí read-modify-write operace nad sdílenou nebo globální pamětí • žádná interference s ostatními vlákny • 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 9 aritmetické (Add, Sub, Exch, Min, Max, lne, Dec, CAS) a bitové (And, Or, Xor) operace Jin Fihpovic Akcelerace výpočtů na GPU oooo Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód ooooooooooo CUDA podrobně OOOOOO0OOO 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 Jin Fihpovic Akcelerace výpočtů na GPU oooo Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód ooooooooooo CUDA podrobně oooooooooo Synchronizace bloků Mezi bloky 9 globální paměť viditelná pro všechny bloky o 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 J in Fihpovic Akcelerace výpočtů na GPU Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód ooooooooooo CUDA podrobně oooooooooo Globální synchronizace pres atomické operace Problém součtu všech prvků vektoru o každý blok sečte prvky své části vektoru • poslední blok sečte výsledky ze všech bloků • implementuje slabší globální bariéru (po dokončení výpočtu u bloků 1..A7 — 1 pokračuje pouze blok n) J in Fihpovic Akcelerace výpočtů na GPU Motivace Architektura GPU C for CUDA Demonstrační kód CUDA podrobně Závěr oooo oooooooooo ooooo ooooooooooo 000000000« oo __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 = atomicInc(&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č Akcelerace výpočtů na GPU 41 /43 Architektura GPU oooooooooo C for CUDA ooooo Demonstrační kód ooooooooooo CUDA podrobně oooooooooo CUDA dokumentace (instalována s CUDA Toolkit, ke stažení na developer, n vidia. com) 9 CUDA C Programming Guide (nejdůležitější vlastnosti CUDA) o CUDA C Best Practices Guide (detailnejší zaměření na optimalizace) • CUDA Reference Manual (kompletní popis C for CUDA API) a další užitečné dokumenty (manuál k nvcc, popis PTX jazyka, manuály knihoven, ...) Série článků CUDA, Supercomputing for the Masses • http://www.ddj.com/cpp/207200659 J in Fihpovic Akcelerace výpočtů na GPU Dnes jsme si ukázali o k čemu je dobré znát CUDA 9 v čem jsou GPU jiná • základy programování v C for CUDA Příště se zaměříme na 9 jak psát efektivní GPU kód 4 ^ >■ < ► 4 -š J in Fihpovic Akcelerace výpočtů na GPU