Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo oooooooooooo ooooo oooooooo 00 Úvod, základy CUDA Jiří Filipovič podzim 2010 Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr •oooooooo ooooooo oooooo oooooooooooo ooooo oooooooo oo Motivace - Moorův zákon Moorův zákon Počet tranzistorů na jednom čipu se přibližně každých 18 měsíců zdvojnásobí. □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr •oooooooo ooooooo oooooo oooooooooooo ooooo oooooooo oo Motivace - Moorův zákon Moorů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 □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr o»ooooooo ooooooo oooooo oooooooooooo ooooo oooooooo oo Motivace - změna paradigmatu Důsledky Moorova zákona: • dříve: rychlost zpracování programového vlákna procesorem se každých 18 měsíců zdvojnásobí • změny ovlivňují především návrh kompilátoru, aplikační programátor se jimi nemusí zabývat • dnes: rychlost zpracování dostatečného počtu programových vláken se každých 18 měsíců zdvojnásobí • pro využití výkonu dnešních procesorů je zapotřebí paralelizovat algoritmy • paralelizace vyžaduje nalezení souběžnosti v řešeném problému, což je (stále) úkol pro programátora, nikoliv kompilátor □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr oo»oooooo ooooooo oooooo oooooooooooo ooooo oooooooo oo Motivace - druhy paralelismu • úlohový paralelismus • problém je dekomponován na úlohy, které mohou být prováděny souběžně » úlohy jsou zpravidla komplexnější, mohou provádět různou činnost • vhodný pro menší počet výkonných jader • zpravidla častější (a složitější) synchronizace • datový paralelismus • souběžnost na úrovni datových struktur • zpravidla prováděna stejná operace nad mnoha prvky datové struktury • jemnější paralelismus umožňuje konstrukci jednodušších procesorů □ g - = = -^c^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooo»ooooo ooooooo oooooo oooooooooooo ooooo oooooooo oo Motivace - druhy paralelismu • z pohledu programátora • rozdílné paradigma znamená rozdílný pohled na návrh algoritmů • některé problémy jsou spíše datově paralelní, některé úlohově • z pohledu vývojáře hardware » procesory pro datově paralelní úlohy mohou být jednodušší • při stejném počtu tranzistorů lze dosáhnout vyššího aritmetického výkonu • jednodušší vzory přístupu do paměti umožňují konstrukci HW s vysokou paměťovou propustností □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr oooo»oooo ooooooo oooooo oooooooooooo ooooo oooooooo oo Motivace - grafické výpočty • datově paralelní • provádíme stejné výpočty pro různé vertexy, pixely, • 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 □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooo«ooo ooooooo oooooo oooooooooooo ooooo oooooooo oo Motivace - výkon Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr oooooo«oo ooooooo oooooo oooooooooooo ooooo oooooooo oo Motivace - výkon Theoretical GB/s 2003 2004 2005 2006 2007 2008 2009 2010 Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooo»o ooooooo oooooo oooooooooooo ooooo oooooooo oo Motivace - shrnutí • GPU jsou výkonné • řádový nárůst výkodu již stoji za studium nového programovacího modelu • pro plné využití moderních GPU i CPU je třeba programovat paralelně • paralelní architektura GPU přestává být řádově náročnější • GPU jsou široce rozšířené • jsou levné • spousta uživatelů má na stole superpočítač □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr 00000000» ooooooo oooooo oooooooooooo ooooo oooooooo oo Motivace - uplatnění Využití GPU pro obecné výpočty je dynamicky se rozvíjející oblast s širokou škálou aplikací □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr 00000000» ooooooo oooooo oooooooooooo ooooo oooooooo oo 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ší... □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr 00000000» ooooooo oooooo oooooooooooo ooooo oooooooo oo 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ódovania dekódování multimediálních dat • herní fyzika • úprava obrázků, 3D rendering • atd... Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr OOOOOOOOO »000000 oooooo oooooooooooo ooooo oooooooo oo Architektura GPU CPU vs. GPU • jednotky jader vs. desítky 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í □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr OOOOOOOOO OfOOOOO oooooo oooooooooooo ooooo oooooooo oo Architektura GPU Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo oosoooo oooooo oooooooooooo ooooo oooooooo oo 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 □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooo»ooo oooooo oooooooooooo ooooo oooooooo oo 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ů a HW přepínání a plánování threadů • thready organizovány po 32 do warpů • SIMT • nativní synchronizace v rámci multiprocesoru □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo oooo»oo oooooo oooooooooooo ooooo oooooooo oo 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 multi procesoru • 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 □ g - = = -^c^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr OOOOOOOOO OOOOO0O oooooo oooooooooooo ooooo oooooooo oo Procesor G80 Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr OOOOOOOOO 000000» oooooo oooooooooooo ooooo oooooooo oo Další vývoj Procesory odvozené od G80 • double-precision výpočty • relaxovány pravidla pro efektivní přístup ke globální paměti • navýšeny on-chip zdroje (více registrů, více threadů na MP) a lepší možnosti synchronizace (atomické operace, hlasování warpů) Fermi • vyšší paralelizace na úrovni multiprocessoru (více jader, dva warp schedulery, více DP výkonu) • konfigurovatelná LI a sdílená L2 cache • plochý adresní prostor » lepší přesnost v plovoucí řádové čárce a paralelní běh kernelů • širší možnosti synchronizace • další změny plynoucí z odlišné architektury Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr OOOOOOOOO OOOOOOO «00000 oooooooooooo ooooo oooooooo oo 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 o je možné použít ji s více programovacími jazyky c OpenCL Fortran C++ MM CUDA Architecture m □ S1 ~ = -š -O Q, o Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo o»oooo oooooooooooo ooooo oooooooo oo 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 Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oo»ooo oooooooooooo ooooo oooooooo oo 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 □ gl - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo ooo«oo oooooooooooo ooooo oooooooo oo Hierarchie vláken Grid Block (O, O) Block (1,0) Block (2, O) Block (O, iy Block (1,1) v Block (2,1) Block (1,1) Thread (0, 0) Thread (1, 0) Thread (2, 0) Thread (3, 0) i Thread (0,1) Thread (1,1) 1 Thread (2,1) 1 Thread (3,1) 1 Thread (0, 2) Thread (1, 2) i Thread (2, 2) i Thread (3, 2) i □ gi - Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooo»o oooooooooooo ooooo oooooooo oo 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 □ gl - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo ooooo« oooooooooooo ooooo oooooooo oo Hierarchie pamětí Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr OOOOOOOOO OOOOOOO OOOOOO »00000000000 ooooo oooooooo oo Příklad - součet vektorů Chceme sečíst vektory a a b a výsledek uložit do vektoru c. Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr OOOOOOOOO OOOOOOO OOOOOO »00000000000 ooooo oooooooo oo 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. □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr OOOOOOOOO OOOOOOO OOOOOO »00000000000 ooooo oooooooo oo 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]; Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr OOOOOOOOO OOOOOOO OOOOOO »00000000000 ooooo oooooooo oo 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. □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr OOOOOOOOO OOOOOOO OOOOOO »00000000000 ooooo oooooooo oo 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ý thread sečte i-té složky vektorů: c[i] = a[i] + b[i]; Jak zjistíme, kolikátý jsme thread? □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo o«oooooooooo ooooo oooooooo oo Hierarchie vláken Grid Block (O, O) Block (1,0) Block (2, O) Block (O, iy Block (1,1) v Block (2,1) Block (1,1) Thread (0, 0) Thread (1, 0) Thread (2, 0) Thread (3, 0) i Thread (0,1) Thread (1,1) 1 Thread (2,1) 1 Thread (3,1) 1 Thread (0, 2) Thread (1, 2) i Thread (2, 2) i Thread (3, 2) i □ @l - Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo oo»ooooooooo ooooo oooooooo oo 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.jx, y, z} udává velikost bloku • blockldx.jx, 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) □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo ooo«oooooooo ooooo oooooooo oo Příklad - součet vektorů Vypočítáme tedy globální pozici threadu (mřížka i bloky jsou jednorozměrné): □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo ooo«oooooooo ooooo oooooooo oo 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; □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo ooo«oooooooo ooooo oooooooo oo 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]; } Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo ooo«oooooooo ooooo oooooooo oo 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. □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo oooo»ooooooo ooooo oooooooo oo 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í □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo ooooo»oooooo ooooo oooooooo 00 Ke kompletnímu výpočtu je třeba: Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo ooooo»oooooo ooooo oooooooo 00 Ke kompletnímu výpočtu je třeba: • alokovat paměť pro vektory, naplnit je daty □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo ooooo»oooooo ooooo oooooooo 00 Ke kompletnímu výpočtu je třeba: • alokovat paměť pro vektory, naplnit je daty • alokovat paměť na GPU □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo ooooo»oooooo ooooo oooooooo oo 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 □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo ooooo»oooooo ooooo oooooooo oo 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 □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo ooooo»oooooo ooooo oooooooo oo 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 □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo ooooo»oooooo ooooo oooooooo oo 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 :-) □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo oooooo»ooooo ooooo oooooooo oo Příklad - součet vektorů CPU kód naplní a a b, vypíše c: #include #define 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č Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo ooooooosoooo ooooo oooooooo oo Správa GPU paměti Paměť je třeba dynamicky alokovat. cudaMal1oc(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 cuda MemcpyDevice ToHosť). Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo oooooooo»ooo ooooo oooooooo oo 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)); cudaMalloc((void**)&d_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); □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo ooooooooosoo ooooo oooooooo oo 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: #define BLOCK 32 addvec«(d_a , d_b , d_c); Jak řešit problém pro obecnou velikost vektoru? □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo oooooooooo»o ooooo oooooooo oo Příklad - součet vektorů Upravíme kód kernel u: __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); □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr OOOOOOOOO OOOOOOO OOOOOO OOOOOOOOOOO* ooooo oooooooo oo 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 hi-end GPU: barracuda.fi.muni.cz, airacuda.fi.muni.cz, účty na přání □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU ooooooooo ooooooo CUDA oooooo Demonstrační kód oooooooooooo Paměťová hierarchie •oooo Synchronizace oooooooo Závěr 00 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) □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU ooooooooo ooooooo CUDA oooooo Demonstrační kód oooooooooooo Paměťová hierarchie •oooo Synchronizace oooooooo Závěr 00 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 nevleze 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) □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU ooooooooo ooooooo CUDA oooooo Demonstrační kód oooooooooooo Paměťová hierarchie o»ooo Synchronizace oooooooo Závěr 00 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č Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo oooooooooooo oo»oo oooooooo oo 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 = ( int *)&array 1 [ 1 2 8] ; short *array3 = (short*)&array2[25 6]; 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«>>(); Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU ooooooooo ooooooo CUDA Demonstrační kód oooooo oooooooooooo Paměťová hierarchie ooo»o Synchronizace oooooooo Závěr 00 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 ..device □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo oooooooooooo oooo» oooooooo oo Ostatní paměti • paměť konstant • texturová paměť • systémová paměť □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr OOOOOOOOO OOOOOOO OOOOOO OOOOOOOOOOOO OOOOO »0000000 oo 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 □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo oooooooooooo ooooo o»oooooo oo 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č Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo oooooooooooo ooooo oo»ooooo oo 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ů. Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo oooooooooooo ooooo ooo»oooo oo 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 syncthreadsQ □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo oooooooooooo ooooo oooo«ooo oo 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 □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo oooooooooooo ooooo ooooo»oo oo 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 □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo oooooooooooo ooooo oooooo«o oo Globální synchronizace přes atomické operace Problém součtu všech prvků vektoru • 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 zkončení výpočtu u bloků 1../7 — 1 pokračuje pouze blok n) □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr OOOOOOOOO OOOOOOO OOOOOO OOOOOOOOOOOO OOOOO 0000000» 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 = atomiclnc (&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č Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo oooooooooooo ooooo oooooooo »o Materiály CUDA dokumentace (instalována s CUDA Toolkit, ke stažení na developer, nvidia.com) • CUDA C Programming Guide (nejdůležitější vlastnosti CUDA) • CUDA C Best Practices Guide (detailnější zaměření na optimalizace) • CUDA Reference Manual (kompletní popis C for CUDA API) • další užitečné dokumenty (manuál k nvcc, popis PTX jazyka, manuály knihoven, ...) Textbook ke kurzům na University of Illinois • dostupný z http://courses.ece.illinois.edu/ece498/al/Syllabus.html Série článků CUDA, Supercomputing for the Masses • http://www.ddj.com/cpp/207200659 Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo oooooooooooo ooooo oooooooo o» Dnes jsme si ukázali • k čemu je dobré znát CUDA • v čem jsou GPU jiná • základy programování v C for CUDA □ - = = -0*3*0 Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo oooooooooooo ooooo oooooooo o» Dnes jsme si ukázali • k čemu je dobré znát CUDA • v čem jsou GPU jiná 9 základy programování v C for CUDA Příště se zaměříme na • jak psát efektivní GPU kód □ - = = -0*3*0 Jiří Filipovič Úvod, základy CUDA Motivace Architektura GPU CUDA Demonstrační kód Paměťová hierarchie Synchronizace Závěr ooooooooo ooooooo oooooo oooooooooooo ooooo oooooooo o* Dnes jsme si ukázali • k čemu je dobré znát CUDA • v čem jsou GPU jiná • základy programování v C for CUDA Příště se zaměříme na • jak psát efektivní GPU kód K samostatné práci • zkuste si přeložit první CUDA program • máte-li chuť, experimentujte s ním! Jiří Filipovič Úvod, základy CUDA