0 předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr ooooo ooooooooooo ooooooo oooooo oooooooooooo Úvod, základy CUDA Jiří Filipovič podzim 2010 Jiří Filipovič Úvod, základy CUDA 0 předmětu •oooo 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 grafických procesorech. □ - = = ^q^o Jiří Filipovič Úvod, základy CUDA 0 předmětu •oooo 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 grafických procesorech. Budeme probírat: • návrh paralelních algoritmů s důrazem na užití v programovacím modelu poskytovaném dnešními GPU □ - = = ^q^o Jiří Filipovič Úvod, základy CUDA 0 předmětu •oooo 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 grafických procesorech. Budeme probírat: • návrh paralelních algoritmů s důrazem na užití v programovacím modelu poskytovaném dnešními GPU • architekturu GPU založených na CUDA □ - = = ^q^o Jiří Filipovič Úvod, základy CUDA 0 předmětu •oooo 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 grafických procesorech. Budeme probírat: • návrh paralelních algoritmů s důrazem na užití v programovacím modelu poskytovaném dnešními GPU • architekturu GPU založených na CUDA • programování v C for CUDA □ - = = ^q^o Jiří Filipovič Úvod, základy CUDA 0 předmětu •oooo 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 grafických procesorech. Budeme probírat: • návrh paralelních algoritmů s důrazem na užití v programovacím modelu poskytovaném dnešními GPU • architekturu GPU založených na CUDA • programování v C for CUDA • nástroje a knihovny □ - = = ^q^o Jiří Filipovič Úvod, základy CUDA 0 předmětu •oooo 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 grafických procesorech. Budeme probírat: • návrh paralelních algoritmů s důrazem na užití v programovacím modelu poskytovaném dnešními GPU • architekturu GPU založených na CUDA • programování v C for CUDA • nástroje a knihovny • optimalizaci kódu pro GPU □ - = = ^q^o Jiří Filipovič Úvod, základy CUDA 0 předmětu •oooo 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 grafických procesorech. Budeme probírat: • návrh paralelních algoritmů s důrazem na užití v programovacím modelu poskytovaném dnešními GPU • architekturu GPU založených na CUDA • programování v C for CUDA • nástroje a knihovny • optimalizaci kódu pro GPU • případové studie □ - = = ^q^o Jiří Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód •oooo 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 grafických procesorech. Budeme probírat: • návrh paralelních algoritmů s důrazem na užití v programovacím modelu poskytovaném dnešními GPU • architekturu GPU založených na CUDA • programování v C for CUDA • nástroje a knihovny • optimalizaci kódu pro GPU • případové studie Předmět je prakticky orientován - GPU je konstantně-krát rychlejší než CPU, vedle časové složitosti nás tedy zajímá, jak napsat optimální kód. Jiří Filipovič Úvod, základy CUDA 0 předmětu o»ooo Motivace Architektura GPU ooooooooooo ooooooo CUDA oooooo Demonstrační kód oooooooooooo Závěr 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ů □ - = = ^q^o Jiří Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód oo»oo ooooooooooo ooooooo oooooo oooooooooooo Hodnocení Zakončení zkouškou o A: 92-100 o B: 86-91 o C: 78-85 o D: 72-77 o E: 66-71 o F: 0 - 65 bodů Zakončení kolokviem » 50 bodů □ - = = ^q^o Jiří Filipovič Úvod, základy CUDA 0 předmětu ooo»o 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.illinois.edu/ece498/al/Syllabus.html Série článků CUDA, Supercomputing for the Masses • http://www.ddj.com/cpp/207200659 □ S - = = ^<^o Jiří Filipovič Úvod, základy CUDA 0 předmětu oooo» Motivace Architektura GPU CUDA ooooooooooo ooooooo oooooo Demonstrační kód oooooooooooo Závěr 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 □ - = = ^q^o Jiří Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód OOOOO »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í. □ - = = ^q^o Jiří Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód OOOOO »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 Jiří Filipovič Úvod, základy CUDA 0 předmětu ooooo 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 □ g - = = ^c^o Jiří Filipovič Úvod, základy CUDA 0 předmětu ooooo 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ů □ g - = = ^q^O Jiří Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód ooooo ooo»ooooooo ooooooo oooooo oooooooooooo 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 O předmětu Motivace Architektura GPU CUDA Demonstrační kód ooooo oooo«oooooo ooooooo oooooo oooooooooooo 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 0 předmětu ooooo Motivace Architektura GPU OOOOO0OOOOO ooooooo CUDA oooooo Demonstrační kód oooooooooooo Závěr Motivace - výkon Theoretical GFLOP/s 1500 GeForce GTX 480 Jiří Filipovič Úvod, základy CUDA 0 předmětu ooooc Motivace Architektura GPU oooooo«oooo ooooooo CUDA oooooo Demonstrační kód oooooooooooo Závěr Motivace - výkon Theoretical GB/s 2003 2004 2005 2006 2007 2008 2009 2010 O předmětu Motivace Architektura GPU CUDA Demonstrační kód ooooo ooooooo»ooo ooooooo oooooo oooooooooooo 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 0 předmětu ooooo Motivace Architektura GPU oooooooo»oo ooooooo CUDA oooooo Demonstrační kód oooooooooooo Závěr Motivace - uplatnění Využití GPU pro obecné výpočty je dynamicky se rozvíjející oblast s širokou škálou aplikací □ - = = -0*3*0 Jiří Filipovič Úvod, základy CUDA 0 předmětu ooooo Motivace Architektura GPU oooooooo»oo ooooooo CUDA oooooo Demonstrační kód oooooooooooo Závěr 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 0 předmětu ooooo Motivace Architektura GPU oooooooo»oo ooooooo CUDA oooooo Demonstrační kód oooooooooooo Závěr 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 0 předmětu ooooo 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 ooooo ooooooooo»o ooooooo oooooo oooooooooooo Motivace - uplatnění Vývojáři SW jsou stále ještě nedostatkovým zbožím... Výovjáři schopní psát paralelní SW jsou extrémně nedostatkovým zbožím! □ - = = ^q^o Jiří Filipovič Úvod, základy CUDA 0 předmětu ooooo Motivace Architektura GPU ooooooooo»o ooooooo CUDA oooooo Demonstrační kód oooooooooooo Závěr Motivace - uplatnění Vývojáři SW jsou 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 :-) □ - = = ^q^o Jiří Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód ooooo oooooooooo* ooooooo oooooo oooooooooooo Historická exkurze • SI M D model od 60. let • projekt Solomon u firmy Westinghouse na začátku 60. let • 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 • 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 OOOOO OOOOOOOOOOO »000000 oooooo oooooooooooo 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 Architektura GPU Demonstrační kód Architektura GPU Jiří Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód ooooo ooooooooooo oo»oooo 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 □ g - = = -0*3.0 Jiří Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód ooooo 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ů 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 O předmětu Motivace Architektura GPU CUDA Demonstrační kód ooooo ooooooooooo oooo»oo 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 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 - = = ^q^O Jiří Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr OOOOO OOOOOOOOOOO OOOOO0O oooooo oooooooooooo Procesor G80 Jiří Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód OOOOO OOOOOOOOOOO 000000» oooooo oooooooooooo 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 O předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr OOOOO OOOOOOOOOOO OOOOOOO »00000 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 o je možné použít ji s více programovacími jazyky c OpenCL Fortran C++ MM CUDA Architecture m Jiří Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód ooooo ooooooooooo ooooooo o»oooo oooooooooooo 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 O předmětu Motivace Architektura GPU CUDA Demonstrační kód ooooo 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 □ gl - = = ^Q^O Jiří Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr ooooo ooooooooooo ooooooo ooo«oo oooooooooooo 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 O předmětu Motivace Architektura GPU CUDA Demonstrační kód ooooo ooooooooooo ooooooo oooo»o 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 □ g - = = -0*3.0 Jiří Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód ooooo ooooooooooo ooooooo ooooo« oooooooooooo Hierarchie pamětí Jiří Filipovič Úvod, základy CUDA 0 předmětu ooooo 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. Jiří Filipovič Úvod, základy CUDA 0 předmětu ooooo 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. □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód OOOOO 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]; Jiří Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód OOOOO 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. □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA 0 předmětu ooooo 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, 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 O předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr ooooo ooooooooooo ooooooo oooooo o-»oooooooooo 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 0 předmětu Motivace ooooo ooooooooooo Architektura GPU ooooooo CUDA oooooo Demonstrační kód oosooooooooo Závěr 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 0 předmětu ooooo 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 threadů (mřížka i bloky jsou jednorozměrné): □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA 0 předmětu ooooo 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; □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód ooooo 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 ooooo 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. Jiří Filipovič Úvod, základy CUDA O předmětu Motivace Architektura GPU CUDA Demonstrační kód ooooo 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í □ g - = = -0*3.0 Jiří Filipovič Úvod, základy CUDA 0 předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr ooooo ooooooooooo ooooooo oooooo ooooo»oooooo Ke kompletnímu výpočtu je třeba: Jiří Filipovič Úvod, základy CUDA 0 předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr ooooo ooooooooooo ooooooo oooooo ooooo»oooooo Ke kompletnímu výpočtu je třeba: • alokovat paměť pro vektory, naplnit je daty □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA 0 předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr ooooo 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 □ g - = = ^c^o Jiří Filipovič Úvod, základy CUDA 0 předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr ooooo 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 □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA 0 předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr ooooo 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 □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA 0 předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr ooooo 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 □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA 0 předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr ooooo 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 :-) □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA 0 předmětu ooooo 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: #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 O předmětu Motivace Architektura GPU CUDA Demonstrační kód ooooo ooooooooooo ooooooo oooooo ooooooo»oooo 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ť). 0 předmětu ooooo 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)); 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 O předmětu Motivace Architektura GPU CUDA Demonstrační kód ooooo 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: #define BLOCK 32 addvec«(d_a , d_b , d_c); Jak řešit problém pro obecnou velikost vektoru? □ g - = = -0*3.0 Jiří Filipovič Úvod, základy CUDA 0 předmětu ooooo 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 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); □ g - = = -0*3.0 Jiří Filipovič Úvod, základy CUDA 0 předmětu ooooo Motivace Architektura GPU ooooooooooo ooooooo CUDA oooooo Demonstrační kód ooooooooooo* Závěr 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? o 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 □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA 0 předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr ooooo 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 □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA 0 předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr ooooo 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 □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA 0 předmětu Motivace Architektura GPU CUDA Demonstrační kód Závěr ooooo 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! □ - = = ^Q^O Jiří Filipovič Úvod, základy CUDA