0 předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr ooooooo ooooooooooo OOOOOOO oooooo oooooooooooo Úvod, základy CUDA Jiří Filipovič podzim 2011 Jih Filipovič Úvod, základy CUDA 0 předmětu #000000 Motivace Architektura GPU OOOOOOOOOOO OOOOOOO CUDA OOOOOO Demonstrační kód OOOOOOOOOOOO Závěr Co se budeme učit? Předmět se zabývá návrhem algoritmů a programováním obecných výpočtů na vektorových multiprocesorech. Jih Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr ♦OOOOOO OOOOOOOOOOO OOOOOOO OOOOOO OOOOOOOOOOOO Co se budeme učit? Předmět se zabývá návrhem algoritmů a programováním obecných výpočtů na vektorových multiprocesorech. Primárně se zaměříme na CUDA GPU • GPU jsou zajímavé z hlediska rychlosti a rozšířenosti • C for CUDA vhodné k výuce (jednoduché API, množství materiálů, zdrojových kódů, odladěnost prostředí) • omezení na GPU od nVidie a x86 CPU Jih Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr ♦OOOOOO OOOOOOOOOOO OOOOOOO OOOOOO OOOOOOOOOOOO Co se budeme učit? Předmět se zabývá návrhem algoritmů a programováním obecných výpočtů na vektorových multiprocesorech. Primárně se zaměříme na CUDA GPU • GPU jsou zajímavé z hlediska rychlosti a rozšířenosti • C for CUDA vhodné k výuce (jednoduché API, množství materiálů, zdrojových kódů, odladěnost prostředí) • omezení na GPU od nVidie a x86 CPU Ukážeme si také základy OpenCL • z CUDA poměrně snadný přechod • kompilovatelné pro nVidia a AMD GPU, x86 CPU i Cell • probereme specifika optimalizace pro AMD GPU a x86 CPU Jih Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr ♦OOOOOO OOOOOOOOOOO OOOOOOO OOOOOO OOOOOOOOOOOO Co se budeme učit? Předmět se zabývá návrhem algoritmů a programováním obecných výpočtů na vektorových multiprocesorech. Primárně se zaměříme na CUDA GPU • GPU jsou zajímavé z hlediska rychlosti a rozšířenosti • C for CUDA vhodné k výuce (jednoduché API, množství materiálů, zdrojových kódů, odladěnost prostředí) • omezení na GPU od nVidie a x86 CPU Ukážeme si také základy OpenCL • z CUDA poměrně snadný přechod • kompilovatelné pro nVidia a AMD GPU, x86 CPU i Cell • probereme specifika optimalizace pro AMD GPU a x86 CPU Předmět je prakticky orientován - paralelní výpočet je na jednom procesoru konstantně-krát rychlejší než sekvenční, vedle časové složitosti nás tedy zajímá, jak napsat efektivní^kpqjg, > Jiří Filipovič Úvod, základy CUDA 0 předmětu OÄOOOOO Motivace Architektura GPU OOOOOOOOOOO OOOOOOO CUDA oooooo Demonstrační kód oooooooooooo Závěr Co se budeme učit? Budeme probírat: • návrh datově-paralelních algoritmů s důrazem na užití v GPU Jih Filipovič Úvod, základy CUDA 0 předmětu OÄOOOOO Motivace Architektura GPU OOOOOOOOOOO OOOOOOO CUDA oooooo Demonstrační kód oooooooooooo Závěr Co se budeme učit? Budeme probírat: • návrh datově-paralelních algoritmů s důrazem na užití v GPU • architekturu GPU založených na CUDA Jih Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr 0«00000 OOOOOOOOOOO OOOOOOO OOOOOO OOOOOOOOOOOO Co se budeme učit? Budeme probírat: • návrh datově-paralelních algoritmů s důrazem na užití v GPU • architekturu GPU založených na CUDA • programování v C for CUDA Jih Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr 0«00000 OOOOOOOOOOO OOOOOOO OOOOOO OOOOOOOOOOOO Co se budeme učit? Budeme probírat: • návrh datově-paralelních algoritmů s důrazem na užití v GPU • architekturu GPU založených na CUDA • programování v C for CUDA • nástroje a knihovny Jih Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr 0«00000 OOOOOOOOOOO OOOOOOO OOOOOO OOOOOOOOOOOO Co se budeme učit? Budeme probírat: • návrh datově-paralelních algoritmů s důrazem na užití v GPU • architekturu GPU založených na CUDA • programování v C for CUDA • nástroje a knihovny • optimalizaci kódu pro GPU Jih Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr 0«00000 OOOOOOOOOOO OOOOOOO OOOOOO OOOOOOOOOOOO Co se budeme učit? Budeme probírat: • návrh datově-paralelních algoritmů s důrazem na užití v GPU • architekturu GPU založených na CUDA • programování v C for CUDA • nástroje a knihovny • optimalizaci kódu pro GPU • přechod k OpenCL Jih Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr 0«00000 OOOOOOOOOOO OOOOOOO OOOOOO OOOOOOOOOOOO Co se budeme učit? Budeme probírat: • návrh datově-paralelních algoritmů s důrazem na užití v GPU • architekturu GPU založených na CUDA • programování v C for CUDA • nástroje a knihovny • optimalizaci kódu pro GPU • přechod k OpenCL • specifika optimalizace pro AMD Evergreen a Intel x86 Jih Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr 0«00000 OOOOOOOOOOO OOOOOOO OOOOOO OOOOOOOOOOOO Co se budeme učit? Budeme probírat: • návrh datově-paralelních algoritmů s důrazem na užití v GPU • architekturu GPU založených na CUDA • programování v C for CUDA • nástroje a knihovny • optimalizaci kódu pro GPU • přechod k OpenCL • specifika optimalizace pro AMD Evergreen a Intel x86 • případové studie Jih Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód 00*0000 OOOOOOOOOOO OOOOOOO OOOOOO OOOOOOOOOOOO Co od vás budeme vyžadovat Během semestru budete vypracovávat prakticky orientovaný projekt • významný podíl na výsledném hodnocení • jednotné zadání, budeme srovnávat rychlost implementací • celkem 50 + 30 bodů z výsledné známky • funkční kód: 25 bodů • efektivní implementace: 25 bodů • umístění ve srovnání se spolužáky: 30 bodů (pouze ke zlepšení známky) Zkouška (ústní či písemná dle počtu studentů) • 50 bodů Jih Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód Závět 000*000 ooooooooooo ooooooo oooooo oooooooooooo Hodnocení Zakončení zkouškou 0 A: 92-100 0 B: 86-91 0 C: 78-85 0 D: 72-77 0 E: 66-71 o F: 0 - 65 bodů Zakončení kolokviem • 50 bodů Jih Filipovič Úvod, základy CUDA 0 předmětu 0000*00 Motivace Architektura GPU OOOOOOOOOOO OOOOOOO CUDA oooooo Demonstrační kód oooooooooooo Závěr Materiály - CUDA 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. i llinois.edu/ece498/al/Syl la bus. html Série článků CUDA, Supercomputing for the Masses • http: //www.ddj.com/cpp/207200659 Jih Filipovič Úvod. základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód OOOOO0O OOOOOOOOOOO OOOOOOO OOOOOO OOOOOOOOOOOO Materiály - OpenCL • OpenCL 1.1 Specification • AMD Accelerated Parallel Processing Programming Guide • Intel OpenCL SDK Programming Guide • Writing Optimal OpenCL Code with Intel OpenCL SDK Jih Filipovič Úvod. základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód 000000« ooooooooooo ooooooo oooooo oooooooooooo Materiály - paralelní programování • Ben-Ari M., Principles of Concurrent and Distributed Programming, 2nd Ed. Addison-Wesley, 2006 • Timothy G. Mattson, Beverly A. Sanders, Berna L. Massingill, Patterns for Parallel Programming, Addison-Wesley, 2004 J i ŕ í Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód OOOOOOO »0000000000 ooooooo oooooo oooooooooooo 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í. Jiří Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód OOOOOOO »0000000000 ooooooo oooooo oooooooooooo 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 Jih Filipovič Úvod, základy CUDA 0 předmětu ooooooo Motivace o«ooooooooo Architektura GPU CUDA OOOOOOO oooooo Demonstrační kód OOOOOOOOOOOO Závěr 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 Jih Filipovič Úvod, základy CUDA 0 předmětu ooooooo Motivace Architektura GPU oo»oooooooo ooooooo CUDA oooooo Demonstrační kód oooooooooooo Závěr 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ů Jih Filipovič Úvod, základy CUDA 0 předmětu ooooooo Motivace Architektura GPU ooo«ooooooo ooooooo CUDA oooooo Demonstrační kód oooooooooooo Závěr Motivace - druhy paralelismu Z pohledu programátora • rozdílné paradigma znamená rozdílný pohled na návrh algoritmů • některé problémy jsou spise 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í Jiří Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód ooooooo oooo»oooooo ooooooo oooooo oooooooooooo Motivace - výkon GPU Jih Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód ooooooo ooooo«ooooo ooooooo oooooo oooooooooooo Motivace - výkon GPU Jiří Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód ooooooo oooooo»oooo ooooooo oooooo oooooooooooo Motivace - výkon CPU Dnešní CPU mají více jader a vektorové instrukce. Intel Core2 Quad • 4 jádra, SSE instrukce (128-bitové vektory) • sekvenční výpočet 16x pomalejší než vektorový vícevláknový Intel Ivy Bridge • 4-8(?) jader, AVX instrukce (256-bitové vektory) • sekvenční výpočet až 64x pomalejší než vektorový vícevláknový Jiří Filipovič Úvod. základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód Závět ooooooo ooooooo«ooo ooooooo oooooo oooooooooooo Motivace - shrnutí GPU jsou výkonné • řádový nárůst výkonu již stoji za studium nového programovacího modelu Pro plné využití moderních GPU i CPU je třeba programovat paralelně • desetinásobky rychlosti při plném využití CPU, stonásobky při akceleraci na GPU • paralelizace pro GPU je obecně podobně náročná jako paralelizace a vektorizace pro CPU GPU jsou široce rozšířené • jsou levné • spousta uživatelů má na stole superpočítač Jiří Filipovič Úvod, základy CUDA 0 předmětu OOOOOOO Motivace Architektura GPU 00000000*00 OOOOOOO CUDA oooooo Demonstrační kód oooooooooooo Závěr Motivace - uplatnění Využití akcelerace je dynamicky se rozvíjející oblast s širokou škálou aplikací JiFÍ Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód ooooooo oooooooo«oo ooooooo oooooo oooooooooooo Motivace - uplatnění Využití akcelerace 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ší... Jiří Filipovič Úvod. základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód ooooooo oooooooo«oo ooooooo oooooo oooooooooooo Motivace - uplatnění Využití akcelerace 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č Úvod. základy CUDA 0 předmětu OOOOOOO Motivace Architektura GPU OOOOOOOOOÄO OOOOOOO CUDA oooooo Demonstrační kód oooooooooooo Závěr Motivace - uplatnění Vývojáři SWjsou stále ještě nedostatkovým zbožím... Jiří Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód ooooooo ooooooooo«o ooooooo oooooo oooooooooooo Motivace - uplatnění Vývojáři SWjsou stále ještě nedostatkovým zbožím... Výovjáři schopní psát paralelní SW jsou extrémně nedostatkovým zbožím! Jiří Filipovič Úvod. základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód ooooooo ooooooooo«o ooooooo oooooo oooooooooooo Motivace - uplatnění Vývojáři SWjsou stále ještě nedostatkovým zbožím... Výovjáři schopní psát paralelní SW jsou extrémně nedostatkovým zbožím! Obrovské množství existujícího software není paralelní. • pro zajištění růstu výkonu je jej třeba paralelizovat • někdo to musí udělat :-) Jiří Filipovič Úvod. základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód Závět OOOOOOO OOOOOOOOOO* ooooooo oooooo oooooooooooo Historická exkurze • SIMD model od 60. let • projekt Solomon u firmy Westinghouse na začátku 60. let o posléze předán na University of lllinios jako ILLIAC IV • separátní ALU pro každý datový element - masivně paralelní • původní plán: 256 ALUs, 1 GFLOPS • dokončen v r. 1972, 64 ALUs, 100-150 MFLOPS • v 80.-90. letech: běžné vektorové superpočítače, TOP500 a v dnešních CPU: SSE (x86), ActiVec (PowerPC) • Cg: programování vertex a pixel shaderů na grafických kartách (cca 2003) • CUDA: obecné programování GPU, SIMT model (uvolněno poprvé 15. února 2007) • budoucnost?? • OpenCL • vyšší programovací jazyky, automatická paralelizace Jiří Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód OOOOOOO OOOOOOOOOOO #000000 OOOOOO OOOOOOOOOOOO 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í Jih Filipovič Úvod. základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód OOOOOOO OOOOOOOOOOO o«ooooo OOOOOO OOOOOOOOOOOO Architektura GPU Jih Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód OOOOOOO OOOOOOOOOOO 00*0000 OOOOOO OOOOOOOOOOOO 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 Jih Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód OOOOOOO OOOOOOOOOOO OOOÄOOO OOOOOO OOOOOOOOOOOO 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 Jih Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód ooooooo ooooooooooo oooocoo oooooo oooooooooooo 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č Úvod, základy CUDA 0 předmětu OOOOOOO Motivace Architektura GPU OOOOOOOOOOO ooooo»o CUDA oooooo Demonstrační kód oooooooooooo Závěr Procesor G80 Device Multiprocessor N Multiprocessor 2 Multiprocessor 1 Jih Filipovič Úvod, základy CUDA 0 předmětu Motivace Architektura GPU OOOOOOO ooooooooooo oooooo* CUDA oooooo Demonstrační kód oooooooooooo Závěr 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) • 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 • paralelní běh kernelů • širší možnosti synchronizace • další změny plynoucí z odlišné architektury^ Jiří Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód OOOOOOO OOOOOOOOOOO OOOOOOO •OOOOO OOOOOOOOOOOO 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++ DX11 Compute CUDA Architecture i -00,0 Jiří Filipovič Úvod, základy CUDA 0 předmětu Motivace Architektura GPU OOOOOOO OOOOOOOOOOO OOOOOOO CUDA ocoooo Demonstrační kód oooooooooooo Závěr 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 Jih Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód Závět ooooooo ooooooooooo ooooooo oo«ooo oooooooooooo 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č Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr OOOOOOO ooooooooooo OOOOOOO ooo«oo oooooooooooo 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č Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód OOOOOOO OOOOOOOOOOO OOOOOOO 0000*0 oooooooooooo 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 Jiří Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód ooooooo ooooooooooo ooooooo ooooo* oooooooooooo Hierarchie pamětí Jiří Filipovič Úvod, základy CUDA 0 předmětu OOOOOOO Motivace Architektura GPU OOOOOOOOOOO OOOOOOO CUDA oooooo Demonstrační kód •OOOOOOOOOOO Závěr Příklad - součet vektorů Chceme sečíst vektory a a b a výsledek uložit do vektoru c. Jih Filipovič Úvod, základy CUDA 0 předmětu OOOOOOO Motivace Architektura GPU OOOOOOOOOOO OOOOOOO CUDA oooooo Demonstrační kód •OOOOOOOOOOO Závěr 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. Jih Filipovič Úvod. základy CUDA 0 předmětu OOOOOOO Motivace Architektura GPU OOOOOOOOOOO OOOOOOO CUDA oooooo Demonstrační kód •OOOOOOOOOOO Závěr 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]; Jih Filipovič Úvod. základy CUDA 0 předmětu OOOOOOO Motivace Architektura GPU OOOOOOOOOOO OOOOOOO CUDA oooooo Demonstrační kód •OOOOOOOOOOO Závěr 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. Jih Filipovič Úvod. základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód Závět OOOOOOO OOOOOOOOOOO OOOOOOO OOOOOO «00000000000 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? Jiří Filipovič Úvod. základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr OOOOOOO OOOOOOOOOOO OOOOOOO oooooo o»oooooooooo 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č Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód Závět ooooooo ooooooooooo ooooooo oooooo oo«ooooooooo 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) Jiří Filipovič Úvod, základy CUDA 0 předmětu OOOOOOO Motivace Architektura GPU OOOOOOOOOOO OOOOOOO CUDA oooooo Demonstrační kód ooo»oooooooo Závěr Příklad - součet vektorů Vypočítáme tedy globální pozici threadu (mřížka i bloky jsou jednorozměrné): Jih Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód Závět OOOOOOO OOOOOOOOOOO OOOOOOO oooooo ooo»oooooooo Příklad - součet vektorů Vypočítáme tedy globální pozici threadů (mřížka i bloky jsou jednorozměrné): int i = blockldx.x*blockDim.x + threadldx.x; Jiří Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód Závět ooooooo ooooooooooo ooooooo oooooo ooo»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]; } Jiří Filipovič Úvod. základy CUDA 0 předmětu OOOOOOO Motivace Architektura GPU OOOOOOOOOOO OOOOOOO CUDA oooooo Demonstrační kód ooo»oooooooo Závěr 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. Jih Filipovič Úvod. základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód Závět ooooooo ooooooooooo ooooooo oooooo oooo«ooooooo 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č Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód OOOOOOO OOOOOOOOOOO OOOOOOO OOOOOO OOOOO0OOOOOO Ke kompletnímu výpočtu je třeba: Jih Filipovič Úvod, základy CUDA 0 předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr OOOOOOO OOOOOOOOOOO OOOOOOO oooooo ooooo»oooooo Ke kompletnímu výpočtu je třeba: • alokovat paměť pro vektory, naplnit je daty Jiří Filipovič Úvod, základy CUDA 0 předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr OOOOOOO OOOOOOOOOOO OOOOOOO oooooo ooooo»oooooo Ke kompletnímu výpočtu je třeba: • alokovat paměť pro vektory, naplnit je daty • alokovat paměť na GPU Jiří Filipovič Úvod, základy CUDA 0 předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr OOOOOOO ooooooooooo OOOOOOO oooooo ooooo»oooooo 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č Úvod, základy CUDA 0 předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr ooooooo ooooooooooo OOOOOOO oooooo ooooo»oooooo 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 Jiří Filipovič Úvod, základy CUDA 0 předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr ooooooo ooooooooooo OOOOOOO oooooo ooooo»oooooo 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 Jiří Filipovič Úvod, základy CUDA 0 předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr ooooooo ooooooooooo OOOOOOO oooooo ooooo»oooooo 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 :-) Jiří Filipovič Úvod, základy CUDA 0 předmětu OOOOOOO Motivace Architektura GPU OOOOOOOOOOO OOOOOOO CUDA oooooo Demonstrační kód oooooo«ooooo Závěr 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++) = = 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; } Jih Filipovič Úvod, základy CUDA 0 předmětu Motivace Architektura GPU OOOOOOO OOOOOOOOOOO OOOOOOO CUDA oooooo Demonstrační kód ooooooo»oooo Závěr 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). Jih Filipovič Úvod. základy CUDA 0 předmětu OOOOOOO Motivace Architektura GPU OOOOOOOOOOO OOOOOOO CUDA oooooo Demonstrační kód oooooooo«ooo Závěr 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); Jih Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód Závět ooooooo ooooooooooo ooooooo oooooo ooooooooo»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: Sdefine BLOCK 32 addvec«(d_a , d_b , d_c ) ; Jak řešit problém pro obecnou velikost vektoru? Jiří Filipovič Úvod, základy CUDA 0 předmětu OOOOOOO Motivace Architektura GPU OOOOOOOOOOO OOOOOOO CUDA oooooo Demonstrační kód oooooooooo«o Závěr 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); Jih Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód OOOOOOO OOOOOOOOOOO OOOOOOO OOOOOO 00000000000« 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? • ke vzdálenému připojení: barracuda.fi.muni.cz, airacuda.fi.muni.cz, účty budou vytvořeny • windowsí stanice v učebnách (bude upřesněno) • vlastní stroj: stáhněte a nainstalujte CUDA toolkit a SDK z developer.nvidia.com • zdrojové kódy probírané v přednáškách budeme vystavovat ve studijních materiálech Jiří Filipovič Úvod, základy CUDA 0 předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr OOOOOOO ooooooooooo OOOOOOO OOOOOO oooooooooooo 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 Jiří Filipovič Úvod, základy CUDA 0 předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr ooooooo ooooooooooo OOOOOOO oooooo oooooooooooo 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 • podrobnější seznámení s GPU z hadrwarového hlediska • paralelismus, který nám GPU poskytuje • paměti dostupné pro GPU • ukážeme si složitější příklad GPU implementace Jiří Filipovič Úvod, základy CUDA 0 předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr ooooooo ooooooooooo OOOOOOO oooooo oooooooooooo 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 • podrobnější seznámení s GPU z hadrwarového hlediska • paralelismus, který nám GPU poskytuje • paměti dostupné pro GPU • ukážeme si složitější příklad GPU implementace 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