Úvod do technologie CUDA: Hello CUDA!

27. 7. 2009

Doba čtení: 10 minut

Líbí se vám článek?
Podpořte redakci

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).

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)/hosta. 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/projects/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_kernelu.cu> v hlavni_program.cu, vložíme definice_kernelu.cu sem a hlavni_program.cu do CUFILES)</definice_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/linux/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/linux/emurelease.

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 nvcc₂.0.pdf (referenční příručka k nvcc) a ptx_isa₁.4.pdf (specifikace ptx kódu) dostupné po instalaci CUDA Toolkitu v /usr/local/cude/doc.

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 nvcc₂.0.pdf v /usr/local/cuda/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.