Hlavní navigace

Úvod do technologie CUDA: Hello CUDA!

Lukáš Zaorálek 27. 7. 2009

V minulém díle jsme si vysvětlili co je CUDA, jeho praktické využití a jak nainstalovat CUDA SDK. Zde si i my zkusíme napsat CUDA aplikaci. Než si napíšeme naši “Hello CUDA!” aplikaci, musíme se seznámit s terminologií CUDA. Řekneme si něco k paměťovému modelu a vysvětlíme si pár potřebných API funkcí.

Základní pojmy

Pojďme se tedy podívat na konkrétní jednoduchý příklad pseudo-programu CUDA a na něm si vysvětlit terminologii:

Kód 2.1: jednoduchý kernel a jeho volání

// definice kernelu
__global__ void VecAdd(float* A, float* B, float* C)
{
        //identifikacni cislo vlakna
        int i = threadIdx.x;
        C[i] = A[i] + B[i];
}
int main()
{
        // volani kernelu z funkce main
        VecAdd<<<1, N>>>(A, B, C);
}

Základním stavebním kamenem každé CUDA aplikace je tzv. kernel. Jedná se o funkci, která je definována klíčovým slovem __global__. Tato funkce, tedy kernel, je spouštěn N-krát v jednotlivých vláknech (threads). Každé vlákno, které zpracovává instrukce kernelu, má své identifikační číslo, tzv. thread id. Thread id je programátorovi přístupné přes předdefinovanou proměnnou threadIdx. ThreadIdx je typu dim3, což je 3-složkový (x, y, z) vektor. Možná se nyní ptáte, proč musí být identifikace vlákna definovaná zrovna 3-složkovým vektorem? Je tomu tak proto, že usnadňuje práci s vektory, maticemi, vícedimenzionálními poli. Na druhou stranu je třeba říci, že není nutné a nijak závazné využít všechny tři složky vektoru proměnné threadIdx. Klidně můžeme použít pouze x-složku (jak je znázorněno v pseudokódu 2.1), kde je pouze jednodimenzionální pole. Nyní nastal čas si trošku osvětlit ty podivuhodně vypadající ostré závorky <<<1, N>>> při volání kernelu VecAdd, které rozhodně do standartní syntaxe jazyka C nepatří. Je to specialita syntaxe CUDA programů. K objasnění této záhadně vypadající syntaxe nám bude nápomocen obrázek 2.1 (viz. níže). První číslo uvnitř těchto závorek označuje tzv. mřížku (grid), ve které jsou tzv. bloky (blocks), které jsou v našem příkladu výše označeny proměnnou N, a to opět v ostrých závorkách. Co to znamená? Vysvětlíme nejprve bloky. I když je v našem kódu označená jako proměnná N reprezentující číslo, tak blok není číslo, ale opět se jedná o 3-složkový vektor typu dim3. Na obrázku 2.1 je reprezentována jeho dvoudimenzionální varianta. Opět můžeme, ale nemusíme použít jeho 2. (y) nebo 3. (z) složku. Bloky se pak shlukují do mřížky (grid), která je opět typu dim3. Jinak řečeno: x, y, z-složky mřížky určuje počet bloků v mřížce a x, y, z-složky bloku určují počet vláken v jednotlivých blocích. Pokud je mřížka definována pouze jediným číslem, například 1, bude její x, y, z odpovídat hodnotám 1, 1, 1, pokud bude číslo 4, pak její reprezentace vektorem bude vypadat takto: 4, 1, 1, a nikoliv 4, 4, 4! (jedná se prostě o zkrácený zápis). Stejně tak to platí i pro bloky. Z obrázku 2.1 je dále patrné, jak jsou umístěny vlákna v bloku, kde čísla v závorkách odpovídají x, y souřadnici (proměnné threadIdx) vlákna v bloku. Dále vidíme, že pokud budeme spouštět kernely, každý z nich musí mít definován velikost mřížky, jinými slovy počet bloků, a velikost bloku, tedy počet vláken v bloku. Hranatým závorkám <<<pocetbloku, pocetvlaken="">>> se pak říká execution configuration (může mít víc parametrů, ale o tom později). Ještě jednou malá rekapitulace: pokud budeme mít <<<200, 256>>>, znamená to, že počet bloků je 200 a v každém bloku je 256 vláken. Ale jak již bylo řečeno, je možné definovat počet bloků a vláken přímo pomocí proměnných dim3. Další podstatnou věcí, kterou již můžeme usoudit z výše napsaného, je fakt, že klíčovým slovem __global__ říkáme, že daná funkce bude spouštěna na GPU (device – zařízení), nikoliv na CPU (host). Otázka zní: Existuje nějaké klíčové slovo, kterým bych explicitně řekl, že danou funkci je možné vykonávat jak na CPU, tak na GPU? Ano! Není to jedno klíčové slovo, ale rovnou dvě, __global__ __host__ (ale to trošku předbíhám).

CUDA - 2.1

Obr. 2.1 Mřížky, bloky a vlákna

Když už známe základ, jak vypadá execution configuration (v ostrých <<< >>>) a jak se volají kernely, podívejme se na práci s pamětí. V prvním díle seriálu jsme si řekli pár slov o rozdělení paměti na tzv. globální, sdílenou, texturovací, paměť konstant a registry. Nyní si povíme více o globální paměti. Jak již bylo napsáno v prvním díle, jedná se o paměť, která není kešovaná, to znamená, že přístup k ní je velmi pomalý. Proč ji musíme používat? Protože kopírujeme data z/do paměti zařízení(GPU)/hos­ta. Pokud chceme alokovat paměť pro proměnnou (nebo pole) v globální paměti, zavoláme API funkci cudaMalloc (existují další varianty pro dvourozměrné, resp. třírozměrné pole, tedy cudaMallocPitch, resp. cudaMalloc3D, ale o tom jindy). Obdobně pro uvolnění prostředků globální paměti zavoláme API funkci cudaFree. Vstupní a výstupní parametry těchto funkcí si ukážeme a vysvětlíme v příkladu níže. Protože globální paměť je nekešovaná, i její přístup z kernelu není efektivní. Dovolím si trošku předběhnout a povědět vám, že existuje “optimalizace”, které se říká sdružený přístup do paměti (coalesced memory access). Ta zvyšuje rychlost přístupu do globální paměti (více o sdruženém přístupu do paměti v dalším díle). Možná již bylo řečeno hodně teoretických informací, takže je na čase se na vše podívat prakticky.

Náš první program “Hello CUDA!”

Pojďme si ukázat náš první skutečný (přeložitelný) CUDA program a zúročit tak vše, co jsme se doposud naučili. Nejdříve si vytvoříme adresář, kam uložíme soubor se zdrojovými kódy. Adresář se bude jmenovat HelloCuda. Máme-li nainstalovanou CUDA SDK do domovského adresáře a jmenuje-li se NVIDIA_CUDA_SDK, pak bude postup vypadat následovně:

lukas@pocitac:/home/lukas/NVIDIA_CUDA_SDK$ mkdir projects/HelloCuda
lukas@pocitac:/home/lukas/NVIDIA_CUDA_SDK$ gedit projects/HelloCuda/hello_cuda.cu

A můžeme vložit následující zdrojový kód:

Kód 2.2: program hello_cuda.cu

#include <stdio.h>
/////////////////////
//kernel_hello_cuda//
/////////////////////
__global__ void kernel_hello_cuda (size_t num_chars, char* in, char* out) {
//klicove slovo __global__ definuje, ze se jedna o kernel
           //pokud bude id vlakna vetsi nez je pocet znaku v retezci, skonci
           //nepotrebujeme vic vlaken, nez je nezbytne nutne (kopirujeme vcetne 12.znaku, tedy
prazdneho znaku \0)
           if (threadIdx.x > num_chars) {
                      return;
           }
           //kopirujeme obracene retezec !ADUC olleH tak, aby vznikl retezec Hello CUDA!
           //kazde vlakno vezme jeden znak (z retezce in) a ulozi ho na spravne misto (v retezci out)
           out[threadIdx.x] = in[num_chars-threadIdx.x-1];
}
//////////////////////
//hlavni funkce main//
//////////////////////
int main (int argc, char** argv) {
           //alokace pameti pro retezec znaku !ADUC olleH (nesmime zapomenout na 12. znak \0)
           char* hello_cuda_h = (char*)malloc (sizeof(char)*12);
           //kopirovani retezce znaku !ADUC olleH
           strcpy (hello_cuda_h, "!ADUC olleH");
           //retezec znaku pro !ADUC olleH (alokace retezce bude na GPU)
           char* hello_cuda_reverse_d;
           //retezec znaku pro Hello CUDA! (alokace retezce bude na GPU)
           char* hello_cuda_d;
           //velikost retezce znaku v bytech
           int hello_cuda_size = sizeof(char)*(strlen(hello_cuda_h));
           //alokace retezce na GPU (konkretne v global memory) a predani ukazatele
           //prvni parametr: ukazatel na alokovanou pamet
           //druhy parametr: pozadovana velikost pro alokaci pameti
           //vystupni parametr: je enumerator cudaError_t (o nem v dalsich dilech)
           cudaMalloc ((void**) &hello_cuda_reverse_d, hello_cuda_size);
           //alokace retezce na GPU (konkretne v global memory) a predani ukazatele
           cudaMalloc ((void**) &hello_cuda_d, hello_cuda_size);
           //nastaveni prazdneho retezce hello_cuda_d
           cudaMemset (hello_cuda_d, 0, hello_cuda_size);
           //kopirovani retezce z pameti hosta do pameti zarizeni (GPU), konkretne do globalni pameti
GPU
       //prvni parametr: ukazatel cilove pameti
       //druhy paremetr: ukazatel zdrojove pameti
       //treti parametr: velikost kopirovane pameti v bytech
       //ctvrty parametr: enumerator definujici z jakeho druhu pameti do jakeho druhu pameti bude
kopirovano
       //vystupni parametr: je enumerator cudaError_t
       cudaMemcpy (hello_cuda_reverse_d, hello_cuda_h, hello_cuda_size,
cudaMemcpyHostToDevice);
       //volame kernel
       //exekuce instrukci probiha jiz na zarizeni, nikoliv na hostovi!
       kernel_hello_cuda<<<1,15>>> (strlen(hello_cuda_h), hello_cuda_reverse_d, hello_cuda_d);
       //synchronizace vlaken
       //jinymi slovy: cekej dokud nezkonci vsechny vlakna
       cudaThreadSynchronize();
       //kopirovani retezce z pameti zarizeni zpet do pameti hosta
       cudaMemcpy (hello_cuda_h, hello_cuda_d, hello_cuda_size, cudaMemcpyDeviceToHost);
       //nyni je v retezci hello_cuda_h obracena veta Hello CUDA!, kterou vytiskneme
       printf ("%s\n", hello_cuda_h);
       //uvolneni prostredku v pameti zarizeni (GPU)
       //prvni parametr: ukazatel alokovane pameti pomoci funkce cudaMalloc (neda se pouzit pro
pamet alokovanou funkci malloc!)
       //vystupni parametr: je enumerator cudaError_t
       cudaFree (hello_cuda_reverse_d);
       //uvolneni prostredku v pameti zarizeni (GPU)
       cudaFree (hello_cuda_d);
       //uvolneni prostredku v pameti hosta
       free (hello_cuda_h);
       //ukonceni programu
       return 0;
}

Běžně mají soubory zdrojových kódů CUDA programů příponu .cu.

Jen doplním, že pokud si budeme přát kompilovat zdrojové kódy pomocí progamu make, je v adresáři $NVIDIA_CUDA_SDK/pro­jects/template připravena šablona projektu včetně připraveného Makefile, který je nutný editovat a upravit tak, aby odpovídal našemu projektu, tedy:

lukas@pocitac:/home/lukas/NVIDIA_CUDA_SDK$ cp -r projects/template /projects/jmeno_naseho_projektu
lukas@pocitac:/home/lukas/NVIDIA_CUDA_SDK$ gedit projects/jmeno_naseho_projektu/Makefile

kde jednotlivé proměnné znamenají:

  • EXECUTABLE – jméno spustitelného souboru,
  • CUFILES – soubory .cu obsahující C/C++ kód včetně kernelů
  • CU_DEPS – soubory .cu obsahující kernely (např: máme #include <definice_ker­nelu.cu> v hlavni_program­.cu, vložíme definice_kernelu.cu sem a hlavni_program.cu do CUFILES)</defi­nice_kernelu.cu>
  • CCFILES – C/C++ soubory

Tedy v našem příkladě výše úplně vynecháme CU_DEPS, protože nepoužíváme .cu soubory s definicemi kernelů. Dále vynecháme CCFILES a vystačíme si pouze s EXECUTABLE a CUFILES, tedy Makefile bude vypadat takto (zkrácená verze):

# Add source files here
EXECUTABLE            := hello_cuda
# CUDA source files (compiled with cudacc)
CUFILES               := hello_cuda.cu
# CUDA dependency files
CU_DEPS               :=
# C/C++ source files (compiled with gcc / c++)
CCFILES               :=

Po uložení úprav, vytvoření souborů zdrojových kódů a jejich naplnění můžeme používat make. Výstupní spustitelný soubor bude pak v adresáři $NVIDIA_CUDA_SDK/bin/l­inux/release. V případě, že chceme kompilovat daný kód jako emulaci (spouštět binárku aniž bychom spouštěli kód určený pro GPU na zařízení; ale vše pomocí hosta neboli CPU) použijeme make s parametrem emu=1, tedy:

$ make emu=1

V tomto případě se spustitelný soubor bude nacházet v $NVIDIA_CUDA_SDK/bin/l­inux/emurelea­se.

Seznámení s kompilátorem nvcc

Zdrojový kód souboru hello_cuda.cu je nutné zkompilovat do spustitelné podoby. K tomu budeme potřebovat kompilátor nvcc (přesněji compiler driver). Ten má podobné parametry jako gcc. Program jednoduše zkompilujeme takto:

$ nvcc hello_cuda.cu -I../../common/inc -o hello_cuda

Pokud si přejeme emulovaný kód, použijeme:

$ nvcc hello_cuda.cu -I../../common/inc -o hello_cuda -deviceemu

A spustíme jako každý jiný program:

$ ./hello_cuda

Výstup:

$ Hello CUDA!

Jak přesně probíhá kompilace zdrojových kódů? To už nám napovídá obr. 2.2. Program nvcc vezme zdrojový kód v .cu (nebo i .c/.cpp) souboru a za použití dalších aplikací, jako gcc, cl, apod. zkompiluje “standardní” c/c++ zdrojový kód do spustitelné podoby. Kód určený pro zařízení, je přeložen do ptx (Parallel Thread Execution) kódu (podobný jazyku assembler). Tento jazyk je pak kompilován do binárního kódu pro dané zařízení (například pro G8× nebo G9×, přesněji řečeno pro danou verzi compute capability). Společně se spustitelným kódem se zlinkují i knihovny cudart (CUDA runtime library) a cuda (CUDA core library), nutné k běhu aplikace. Více o kompilaci pojednávají dokumenty nvcc2.0.pdf (referenční příručka k nvcc) a ptx_isa1.4.pdf (specifikace ptx kódu) dostupné po instalaci CUDA Toolkitu v /usr/local/cu­de/doc.

Obr 2.2

Obr. 2.2 Kompilace CUDA programů

Závěr

Kompilátor má spousty dalších parametrů, přičemž některé z nich si ukážeme v dalším díle. Pokud máte zájem se již nyní podívat co vše umí, po instalaci CUDA Toolkit by se měl nacházet v již zmiňovaném dokumentu nvcc2.0.pdf v /usr/local/cu­da/doc. V příštím článku si povíme něco o debuggování CUDA aplikací nebo o sdruženém přístupu do paměťi.

Použitá literatura:

  • [1] CUDA 2.2 Programming Guide
  • [2] CUDA 2.2 QuickStart Guide
  • [3] The CUDA Compiler Driver NVCC

Použité obrázky byly převzaty z těchto odkazů:

Našli jste v článku chybu?

27. 7. 2009 23:25

tomas (neregistrovaný)

Pridavam se k ostatnim, ze clanek je temer nepochopitelny pro cloveka neznaleho GPU a CUDA. Autor mel nejprve radne popsat architekturu GPU, z cehoz vyplyne, proc se zabyvame vecmi jako bloky a gridy – je to kvuli synchronizaci mezi vlakny. Bez vylozeni, jak se zpracovavaji vlakna v ramci bloku a gridu nema smysl jit dal a zminovat veci jako coalesced memory acces. Pak bych rekl, co je to kernel a jak se spousti. A potom by zrejme bylo dobre podrobne vylozit, jak dostat data ke zpracovani do glo…

27. 7. 2009 3:52

Ja musim dodat, ze clanek jsem sice necetl, nebyl na to ted vecer cas a nalada, ale projel jsem ho ockem a co se mi nezdalo prehledne, byly zdrojaky, kde se diky jednotne barve michaji komentare s kodem a dost spatne se na to diva. Minimalne jina barva pro komentare by nevadila…

120na80.cz: Pánové, pečujte o svoje přirození a prostatu

Pánové, pečujte o svoje přirození a prostatu

Lupa.cz: Proč firmy málo chrání data? Chovají se logicky

Proč firmy málo chrání data? Chovají se logicky

DigiZone.cz: ČRo rozšiřuje DAB do Berouna

ČRo rozšiřuje DAB do Berouna

Lupa.cz: Propustili je z Avastu, už po nich sahá ESET

Propustili je z Avastu, už po nich sahá ESET

Podnikatel.cz: Přehledná titulka, průvodci, responzivita

Přehledná titulka, průvodci, responzivita

Lupa.cz: Insolvenční řízení kvůli cookies? Vítejte v ČR

Insolvenční řízení kvůli cookies? Vítejte v ČR

Vitalia.cz: Tesco: Chudá rodina si koupí levné polské kuře

Tesco: Chudá rodina si koupí levné polské kuře

Podnikatel.cz: 1. den EET? Problémy s pokladnami

1. den EET? Problémy s pokladnami

Lupa.cz: Avast po spojení s AVG propustí 700 lidí

Avast po spojení s AVG propustí 700 lidí

Podnikatel.cz: Víme první výsledky doby odezvy #EET

Víme první výsledky doby odezvy #EET

120na80.cz: Rakovina oka. Jak ji poznáte?

Rakovina oka. Jak ji poznáte?

Měšec.cz: Finančním poradcům hrozí vracení provizí

Finančním poradcům hrozí vracení provizí

Měšec.cz: Zdravotní a sociální pojištění 2017: Připlatíte

Zdravotní a sociální pojištění 2017: Připlatíte

Podnikatel.cz: Podnikatelům dorazí varování od BSA

Podnikatelům dorazí varování od BSA

DigiZone.cz: Sony KD-55XD8005 s Android 6.0

Sony KD-55XD8005 s Android 6.0

DigiZone.cz: NG natáčí v Praze seriál o Einsteinovi

NG natáčí v Praze seriál o Einsteinovi

Podnikatel.cz: Prodává přes internet. Kdy platí zdravotko?

Prodává přes internet. Kdy platí zdravotko?

Vitalia.cz: Baletky propagují zdravotní superpostel

Baletky propagují zdravotní superpostel

Vitalia.cz: 9 největších mýtů o mase

9 největších mýtů o mase

DigiZone.cz: ČT má dalšího zástupce v EBU

ČT má dalšího zástupce v EBU