Něco málo k úvodu

 

J ak jistě víte, celý počítačový svět jde dopředu a grafické karty nejsou vyjímkou. Stále se kladou větší nároky na výpočetní výkon a rychlost. To vedlo postupně výrobce grafických čipů k několika důležitým změnám v architektuře grafického jádra.Totiž když se díváte na obraz, díváte se vlastně  na čísla, každý pixel je vyjádřený kombinací barev RGB. Krom toho, když kreslíte nějaký 3D objekt, potřebujete například spočítat jeho vzdálenost a polohu. To jsou samozřejmě všechno opět čísla, jen nějak zastíněná v grafickém prostředí. Otázka je tedy jednoduchá: Když se na GPU (Graphics Processing Unit) vyhodnocují číselné operace, nešlo by na něm přímo počítat zcela obecně aniž bychom to musely vnášet do nějakých grafických obálek? Samozřejmě že šlo a CUDA je danou odpovědí. Tento programovací jazyk umožňuje využít celý grafický výkon k ostatním ,,obecným” výpočtům.

Vývojové Prostředí

s26xxz

Pokud pracujete pod Windows, stáhněte si zdarma verzi Microsoft Visual Studia (V současné době,tj 18.12.12 bych doporučil Visual Studio 2010). Visual Studio 2012 sice vypadá o něco lépe, ale instalační balíček CUDY ho nerozpozná a bude vám tvrdit, že ho sice nainstaluje, ale že nevytvoří potřebné šablony pro tvoření nových projektů. Ty jdou samozřejmě vytvořit i ručně a CUDU lze přesvědčit aby běhala i pod 2012, ale obecně se dá říci, že to není vůbec jednoduché. Až Bude Studio nainstalované, jděte na https://developer.nvidia.com/cuda-downloads a stáhněte si dle libosti CUDA 5 nebo CUDA 4.2 (Doporučuji první). Instalační soubor je poněkud veliký, ale měl by obsahovat naprosto vše, tj:

  • Debugger
  • Překladač
  • Nové ovladače
  • Ukázkové kódy
  • Visual Profiler
  • nVidia nSight

Kód napsaný v novém projektu (Například Cuda 5.0 runtime) má koncovku ,, .cu ” což značí Kernel, o kterém se zmíníme při výkladu jazyka. Z hlediska funkce visual studia je ale dobré vědět, že Visual studio přistupuje ke Cuda kódu jako ke klasickému C/C++ . Při překladu se akorát pomocí překladače NVCC přeloží kód, který poběží na Grafické kartě a o zbytek – tj. všechen C/C++ kód – se postará Visual Studio. V projektu můžete vkládat hlavičkové soubory a pracovat s ním zcela jako s normálním C/C++ projektem. Zde bych ale předem rád podotkl, že k Debuggování CUDA kódu je dobré mít a využívat nVidia nSight, který by měl být součástí instalace samotné Cudy. Z hlediska principu si totiž také musíme uvědomit, že adresový prostor na straně CPU má jiné ukazatele, než straně GPU. Tedy nemůžeme použít stejný ukazatel v C/C++ a Cudě.

 

Obecný popis a úvod do jazyka

NV_DesignedFor_CUDA_3D_sm

CUDA se vyznačuje tím, že je to paralelní programovací jazyk. Abychom lépe pochopili slovo paralelní, ukážeme si jednoduchý příklad: Chcete se dívat na libovolnou scénu, která je složená z jednotlivých pixelů, dejme tomu 1920 x 1080. V klasickém sekvenčním prostředí na CPU (Central processing Unit) byste musely vytvořit 2 For cykly a procházet celé dvourozměrné pole a pro každý pixel zvlášť spočítat barvu objektu, jeho polohu, natočení atd … To jsou dva milióny operací ! Naproti tomu CUDA a obecně grafické výpočetní jádro funguje tak, že si vytvoří X*Y vláken (pokud je to možné) a údaje pro všechny pixely najednou …. takže žádné čekání. Jedinou nevýhodou je akorát to, že abyste něco takto spočetli, potřebujete dostat data do grafické paměti (Případně je pak i vykreslit) , ale pokud chcete jen něco spočítat, budete také chtít ty výsledky vrátit zpět – další kopírování.I přes poměrně rychlou sběrnici je to proces časově náročný, takže se obecně nevyplácí v Cudě často kopírovat z Operační paměti RAM na GDDR a zpět.

Samotné spuštění kódu a struktura jazyka je takováto: CUDA používá takzvané Kernely. Můžeme je chápat jako kusy kódu, které provede naprosto každé spuštěné vlákno. Když je kernel spuštěn, musí mu být sděleno, jaká je jeho dimenze (tj kolik dílčích vláken se má spustit. ) To se provádí ve dvou krocích. Nejprve se Kernelu řekne, kolik má spustit bloků a poté se mu řekne, kolik bude mít každý blok vláken. Obojí může být dvourozměrné až třírozměrné číslo. Nicméně to teď není podstatné. Důležité je pochopení funkce bloků a vláken… Teď si asi budete říkat: No dobře, ale jak já poznám, které vlákno má pracovat na jakém úseku? CUDA má opět řešení: každý blok má totiž svůj vlastní index stejně jako každé vlákno. Všechny indexy jsou třírozměrné. Blok navíc disponuje položkou ,,Dimension” což je délka Bloku v daném směru X,Y, nebo Z. Pokud spustíme jen jeden blok, nemusí nás nic moc trápit. Nicméně i v moderních kartách je často udávané omezení, kolik vláken může běžet v rámci jednoho bloku (Obvykle 512-1024 vláken). Takže můžeme počítat s polem o 1000 prvcích.

To je ale poněkud málo, takže se většinou bloků spouští víc. Například pokud spustím 2 bloky a budu mít v každém bloku 512 vláken, mohu počítat s polem o vstupní délce 1024 prvků. Nicméně nemůžu do kernelu napsat, aby se prvky přemapovaly jen jako samotné indexy vláken. V obou blocích totiž sahají indexy vláken jen do 512 a prvky 513 a výše by se nespočetli. Proto je nutné vytvořit si globální index v daném směru (například osy x). Ten získáme snadno, zjistíme si velikost bloku v daném směru a vynásobíme ho indexem bloku, k tomu pak ještě přidáme index vlákna v daném bloku a data vstupního pole jsou namapována. Operace se běžne provádí takto:

 

 

Dá se říct,že čím více vláken běží najednou a v jednom bloku,tím rychlejší je výpočet. Platí to ale jen do určité míry. Pokud je například kód velmi zdlouhavý a potřebuje ke své práci spoustu proměnných, jednotlivé bloky jsou zpomaleny nedostatkem paměti, které má samozřejmě každý blok omezený počet. Abych se vyjádřil přesněji, tak Grafický čip používá několik druhů pamětí: Globální (Ta je většinou udávaná v Gb) … Poté jsou ovšem na kartě v samotném jádru dílčí ,,lokální” paměti a registry. Rozdíl je následovný : do globální paměti může přistupovat každé vlákno, zatímco do lokální (__shared__) paměti mohou přistupovat jen všechna vlákna v rámci jednoho bloku. Registry spadají také pod bloky a slouží k ukládání hodnot v průběhu výpočtu jako lokální proměnné (Typicky například ,,i” ve for cyklu). Co se rychlostí přístupu do těchto pamětí týče, tak platí: Nejrychlejší je přístup do lokální paměti a nejdelší je přístup do globální paměti. Takže se často vyplatí na začátku si data překopírovat z globální do sdílené paměti a pak přistupovat případně do lokální paměti. Nicméně ani to není zcela košér.

 

Základní syntaxe

Spolupráci a základní syntaxi jazyka si ukážeme na jednoduchém příkladu kdy sečteme 2 vstupní pole (budeme tedy sčítat 2 vektory):

 

Na příkladu vidíme použití jednotlivých částí kódů. Nejprve musíme alokovat paměti a to jak pro grafickou paměť, tak klasickou paměť RAM. Rozdíl je v tom, že pro alokování grafické paměti musíme použít příkazu cudaMalloc. Výsledek je uplně stejný jako s funkcí malloc na CPU. V závěru kódu si rovnou můžeme povšimnout, že po použití paměti ji opět uvolníme příkazem cudaFree.  K vysvětlivce HANDLE_ERROR() – Je to dodatečný kód z knížky Cuda By Example. Já osobně jsem si na něj zvykl a stále ho používám, protože je velmi jednoduchý. Má zajistit ukončení aplikace, pokud se v aplikaci vyskytne nějaký error, například pokud ,,šaháme” mimo alokovanou paměť apod. Dále je třeba naplnit vstupní pole nějakými hodnotami, to uděláme na CPU jedním cyklem. Jakmile máme vygenerované nějaké hodnoty,musíme je překopírovat do paměti GPU příkazem cudaMemcpy(),tam uvádíme 4 argumenty: 1 -kam chceme kopírovat. 2- odkud chceme kopírovat. 3- kolik bajtů chceme překopírovat. 4- jaký je směr kopírování: tzn buď DeviceToHost a nebo HostTo Device (Ač to není úplný výčet). V této terminologii je grafická karta označována jako Device a CPU jako Host,z čehož si lze domyslet dané směry kopírování. Zároveň si tedy můžeme povšimnout, že po spuštění kernelu opět výsledky překopírujeme zpět. Nejdůležitější je ale tento kousek kódu:

 

Říká nám,že požadujeme spustit daný kernel s jedním blokem , 128 vlákny na blok a vstupními parametry (ukazately na vstupní data na straně GPU).  Zde bych podotkl, že rovnost: ,,počet bloků * počet vláken na blok    se nemusí rovnat  délce vstupního pole “. Klidně mohu ve stejném příkladu spustit 2 bloky, celý druhý blok ale bude přistupovat do nealokované paměti a aplikace skončí chybou. Tu je třeba ošetřit a dělá se to velmi jednoduše jednou podmínkou:

 

Touto podmínkou zaručíme, že ikdyž budu spuštěno více vláken, daná vlákna nikdy nebudou přistupovat do paměti. Samotný kód kernelu je pak samozřejmě neuvěřitelně složitý: Každé vlákno si přečte jeden prvek z pole ,,a” a jeden prvek z pole ,,b”,pak je sečte a uloží na odpovídající místo do pole ,,C” . Samozřejmě takovýto úkol je velmi triviální, ale slouží k pochopení tohoto programovacího jazyka.

 

CUDA Events

  •  Zjištění časové náročnosti aplikace

Jak už to tak bývá, občas chceme porovnat i některé varianty některého z použitých algoritmů. Můžeme však chtít vědět čas potřebný k provedení daného kódu i jentak čistě pro zajímavost. Následující kód by měl posloužit přesně pro tyto účely. CUDA sama si potřebuje vytvořit 2 proměnné, do které uloží časové otisky na GPU (tj. cudaEventRecord) a poté z jejich rozdílu zjistí čas, jak dlouho daná aplikace běžela (cudaEventElapsedTime). To samo o sobě by ale nestačilo, protože data jsou zpracovávána paralelně a nikde není zaručeno,že ve chvíly sejmutí druhého  otisku už budou data zpracována. Aby tomu tak bylo,musíme zavolat funkci cudaEventSynchrononize,která nám zaručí,že při sejmutí druhého otisku už budou všekna vlákna s prací hotova. Výsledek je pak uváděn v ms.

 

 

Julia Set

Jelikož je sčítání vektorů sice práce zábavná, ale nikoliv zajímavá, ukážeme si další příklad s použitím knihovny OpenGL. Náš příklad znázorní Juliovu množinu pomocí GPU. Množinu přitom vytvoříme tak, že si zvolíme prostor, kde ji budeme hledat a danou komplexní konstantu. Dále si stanovíme počet iterací a kritérium konvergence. Pro každý bod v našem prostoru pak bude platit, že pokud stálým mocněním a přičáním dané konstanty konverguje, bod leží v množině. Pokud diverguje, neleží v množine. Zde je dobré si uvědomit, co je to konvergence a divergence. Tyto matemické pojmy souvisejí s řadami čísel a stručně řečeno se dá říct, že daná řada konverguje, pokud je jejím součtem nějaké reálné číslo, tzn. pokud to není nekonečno.

Jelikož již zdefinice počítáme s komplexními čísly, musíme si zavést strukturu, kde bude definované, jak se má s takovým číslem zacházet. Obecně se tady nebudu rozepisovat o tom, jak funguje mocnění a sčítání komplexních čísel a nepovažuji to za důležité. N začátku kódu si můžeme povšimnout, že jsme si zvolili DIM 1000, to je velikost prostoru v jednom směru. Tedy obrázek bude mít DIM*DIM bodů. V hlavní části programu toho pak není moc zajímavého až na dva řádky:

 První “dim3” je speciální Cuda funkce, která nám umožní vytvářet vícerozměrné proměnné, které pak pustíme do argumentu kernelu. V našem případě chceme pouze 2 rozměry X a Y o velikosti DIM. V následujícím řádku tedy spustíme kernel,který bude mít DIM*DIM bloků a 1 vlákno na blok. Takto vyřešené mapování je výhodné, pokud si to jednak můžeme dovolit a jednak pokud je kód náročný a  jednotlivá vlákna by v rámci jednoho bloku neměla dost výpočetní síly.Takto má každé vlákno blok pro sebe.

Další zajímavou změnou v našem kódu jsou pak doplňky __device__ které říkají Cudě,že daný objekt,nebo funkce bude volána a využívána jen a pouze v Kernelu. Nemůžeme tedy tyto funkce používat v rámci CPU. Co do srovnání s C funkcemi jsou ale téměř naprosto stejné a mají naprosto stejnou strukturu. Naše funkce julia bude počítat,jestli daný bod o souřadnicích x,y leží v dané množině, nebo ne. Na začátku jen posuneme střed osy X a Y na střed, abychom výsledek vůbec viděli.K samotnému výpočtu  jsme si ve funkci definovaly konstantu C, počet iterací (200) a kritérium konvergence (1000). Všechny tyto hodnoty lze měnit, ale  nejzajímavější je asi měnit konstantu.

 

Průběh Kernelu je pak víceméně srozumitelný, nejprve musíme naindexovat vlákna, respektive bloky a každému z nich říct, ať vyhodnotí svůj bod pomocí funkce julia. Vzhledem k tomu, že jako výstupní používame proměnnou typu unsigned char,má každý bod 4 hodnoty. Pomocí jejich kombinace můžeme měnit barvu výsledného zobrazení (Nicméně zelená se myslím hodí báječně).Za zmínku stojí i to, že celý kód samozřejmě můžeme předělat do sekvenční podoby na CPU, kdy budeme ve dvou for cyklech procházet všechny pixely a dělat s nimi úplně to samé. Výsledek celého kódu pak můžeme vidět na obrázku níže,ten sice není přímo z této aplikace,ale je velmi podobný a slouží jako ilustrativní.

Celý kód ve VisualStudiu je rovněž zde: JuliaSet.

 

Problémy s Visual Studiem

  • LINK : fatal error LNK1123: failure during conversion to COFF: file invalid or corrupt

Řešení: Project -> Project Properties -> configuration properties ->  linker -> general -> enable incremental linking -> Set NO.

  • Používáte grafickou knihovnu GL (nebo jeden z naincludovaných souborů)  a pořád vám hlásí hromadu různých chyb.

Řešení: Stahněte si soubor cuda.zip z této stránky (odkaz dole) a do složky C:/windows/system32 si překopírujte soubor glut32.dll  Project->properities->Linker->input->add  “glut32.lib”.  Project->properties->Linker->general->additional input directories>new-> browse to a directory with glut32.lib.

Dále bych doporučil vytvořit si v projektu nové hlavičkkové soubory:

                                        glut.h      glext.h      gl_helper.h

 

Doporučená literatura

Doporučuiji navštívit oficiální vývojářské stránky nVidie na adrese:  developer.nvidia.com . Naleznete tam spoustu rad a návodů na instalaci vývojového prostředí a zároveň odkazy na literaturu, z nichž jednoznačně doporučuji pro začátečníky Knížku CUDA by Example – An Introduction to Parallel Computing. Zároveň pro zájemce uvedu, že 20.12.12 by měla vyjít nové knížka: CUDA Programming – A developers Guide to Parallel Computing. (V současné době již dostupná). Všechny tyto knížky jsou sepsané přímo vývojáři od nVidie a jsou poměrně i čtivé a vyjímečně i úsměvné. Navíc na níže zmíněné stránce CUDA by Example naleznete i zdrojové soubory. V případě, že stránky nebudou dostupné, tak poslední odkaz zahrnuje většinu souborů, které užijete, pokud si také budete chtít postupně procházet jednotlivé příklady.

9780131387683

Závěrem

Cuda toho nabízí mnohem více, toto je jen malý zlomek funkcí. Pokud se chcete naučit využívat jiné druhy paměti, synchronizovat vlákna, koordinovat svůj kód s grafickou knihovnou OpenGl, nebo DirectX, či budete chtít provádět nějaké operace v rámci GPU sekvenčně apod …. Určitě si přečtěte nějakou ze zde uvedených knížek. Oba tyto dva příklady víceméně vycházejí z knížky Cuda By Example. Kdybyste si případně chtěl někdo hrát, můžu uvést další příspěvky, nebo návody, přecijen je to do jisté míry můj koníček, na který ale bohužel nemám tolik času, kolik bych si přál.