4.3 Specializovan´ e knihovny pro pr´ aci s grafick´ ymi kartami
4.3.3 nVidia CUDA
Zvyˇsuj´ıc´ı se poˇzadavky na v´ypoˇcetn´ı v´ykon grafick´ych karet ve 3D aplikac´ıch vedly k tomu, ˇ
ze se bˇeˇzn´e GPU vyvinuly v masivnˇe paraleln´ı, v´ıcevl´aknov´e, v´ıce-j´adrov´e procesory s ob-rovsk´ym v´ypoˇcetn´ım v´ykonem. D˚uvodem takov´ehoto v´yvoje je fakt, ˇze GPU se sv´ymi paraleln´ımi v´ypoˇcty je pˇresnˇe to, co modern´ı real-time grafika potˇrebuje pro dostateˇcnou rychlost vykreslov´an´ı.
2Z hlediska d˚uleˇzitosti architektury CUDA pro tuto pr´aci by bylo vhodn´e tuto ˇc´ast zahrnout jako ka-pitolu na vyˇsˇs´ı ´urovni hierarchie, nicm´enˇe z hlediska logick´eho ˇclenˇen´ı pr´ace i CUDA patˇr´ı do kapitoly o specializovan´ych knihovn´ach pro grafick´e karty, proto se v dalˇs´ım textu budu drˇzet tohoto ˇclenˇen´ı pr´ace.
3Single Instruction Multiple Data
Obr´azek 4.3: Porovn´an´ı CPU a GPU
Obr´azek 4.3 demonstruje rozd´ıly mezi architekturou CPU a GPU, kde v GPU je pouˇzito v´ıcen´asobn´e mnoˇzstv´ı tranzistor˚u pro v´ypoˇcty nad daty, neˇz je tomu u bˇeˇzn´ych CPU. GPU je tedy velmi dobˇre navrˇzeno pro v´ypoˇcty datovˇe paraleln´ıch probl´em˚u. Jedn´a se tedy o pˇr´ıstup, kdy se jeden program prov´ad´ı nad mnoha daty4. Na z´akladˇe v´yˇse uveden´eho tedy vznikla architektura CUDA.
nVidia CUDA (Compute Unified Device Architecture) je relativnˇe nov´a paraleln´ı v´ypoˇcetn´ı architektura vyv´ıjen´a firmou nVidia. CUDA je tedy v´ypoˇcetn´ı n´astroj pro GPU firmy nVi-dia, kter´y je v´yvoj´aˇri pˇr´ıstupn´y pˇres standardn´ı programovac´ı jazyky, jako jsou C ˇci C++
[10] a v souˇcasn´e dobˇe je jiˇz portov´ana do jazyk˚u jako java ˇci Python.
CUDA byla firmou nVidia pˇredstavena v roce 2006 jako paraleln´ı architektura s nov´ym paraleln´ım programovac´ım modelem a instrukˇcn´ı sadou zamˇeˇrenou na paraleln´ı prov´adˇen´ı v´ypoˇct˚u v GPU. V´yvoj CUDA byl podˇr´ızen tomu, aby za cenu relativnˇe mal´eho rozˇs´ıˇren´ı bˇeˇzn´eho programovac´ıho jazyka5 poskytovala jednoduchou a pˇr´ımoˇcarou implementaci pa-raleln´ıch algoritm˚u. CUDA rovnˇeˇz podporuje heterogenn´ı v´ypoˇcetn´ı model, takˇze apli-kace m˚uˇze vyuˇz´ıvat jak v´ypoˇct˚u v CPU, tak v´ypoˇct˚u v GPU. CPU a GPU jsou od sebe vz´ajemnˇe oddˇeleny a chovaj´ı se kaˇzd´e jako samostatn´e programovateln´e zaˇr´ızen´ı a maj´ı sv´e vlastn´ı pamˇet’ov´e modely (prostory). GPU, kter´e je schopno vyuˇz´ıvat CUDA, sest´av´a tedy z mnoha jader, kter´e mohou souˇcasnˇe prov´adˇet mnoho v´ypoˇcetn´ıch vl´aken. Sd´ılen´a pamˇet’
na ˇcipu pak umoˇzˇnuje vl´akn˚um bˇeˇz´ıc´ım paralelnˇe sd´ılet data bez nutnosti jejich pos´ıl´an´ı po syst´emov´e sbˇernici.
V´ypoˇcetn´ı model CUDA
Nyn´ı je vhodn´e se zamˇeˇrit na d˚uleˇzit´e pojmy architektury CUDA6.
• Host
Pojem host v architektuˇre CUDA oznaˇcuje CPU. Zde tedy bˇeˇz´ı program napˇr. v jazyce C.
• Device
Pojemdevice pak oznaˇcuje GPU, v nˇemˇz bude prov´adˇen paraleln´ı program -kernel.
4pˇr´ıstup SIMD - jedna sada instrukc´ı nad mnoha daty
5Napˇr. jazyka C
6V n´asleduj´ıc´ım textu budu pouˇz´ıvat anglick´e term´ıny tak, jak jsou uvedeny v dokumentaci CUDA tam, kde nebyly zavedeny ˇcesk´e ekvivalenty.
Obr´azek 4.4: CUDA - hardwarov´y model, viz [4]
• Kernely (Kernels)
CUDA pˇrin´aˇs´ı program´atorovi moˇznost definovat funkci - kernel, kter´a je po za-vol´an´ı vykon´ana N−kr´at paralelnˇe N r˚uzn´ymi CUDA vl´akny. V z´apisu programu vypad´a jako bˇeˇzn´a funkce jazyka C. Definice kernelu je je uvozena kl´ıˇcov´ym slovem __global__, pot´e zpravidla n´asleduje kl´ıˇcov´e slovovoida n´azev funkce. Poˇcet vl´aken je pak specifikov´an pro kaˇzd´e vol´an´ı kernelu za pouˇzit´ı ”z´avorek” <<<...>>>.
//definice kernelu
__global__ void ScitejMatice(float *a, float *b, float *c){
int i = threadIdx.x;
c[i] = a[i]+b[i];
} ...
int main(int argc, char **argv){
...
Obr´azek 4.5: CUDA - hierarchie vl´aken, viz [4]
ScitejMatice<<<1, N>>>(A, B, C); //vol´an´ı kernelu ...
}
Kernel tedy specifikuje program, kter´y se bude v grafick´e kartˇe prov´adˇet.
• Vl´akno (Thread)
Vl´akno je jakousi nejmenˇs´ı moˇznou jednotkou prov´adˇen´ı v architektuˇre CUDA. Pro tato vl´akna plat´ı v´yˇse uveden´y pˇredpoklad o velmi mal´e aˇz t´emˇeˇr nulov´e reˇzii. Na druhou stranu v´ykon samostatn´eho vl´akna je velmi mal´y, proto, aby byla vyuˇzita potenci´aln´ı v´ypoˇcetn´ı kapacita GPU, je zapotˇreb´ı mnoha (aˇz tis´ıc˚u) vl´aken, kter´a bˇeˇz´ı soubˇeˇznˇe.
Kaˇzd´e vl´akno prov´adˇen´e kernelem m´a svou jednoznaˇcnou identifikaci. Tato identifi-kace je pˇr´ıstupn´a v kernelu skrze vestavˇenou promˇennouthreadIdx, coˇz je zpravidla 3-sloˇzkov´y vektor, takˇze vl´akna mohou b´yt identifikov´ana jedno aˇz tˇr´ı-dimenzion´aln´ım indexem7 vl´akna. Vl´akna mohou b´yt d´ale seskupov´ana do blok˚u vl´aken, kter´e mohou b´yt opˇet aˇz tˇr´ı-dimenzion´aln´ı.
7thread-index
• Warp
Vl´akna jsou v r´amci jednoho bloku seskupena do tzv. wapru. Je to vlastnˇe skupina vl´aken v r´amci bloku, kter´e prov´adˇej´ı stejnou operaci nad r˚uzn´ymi daty. Kaˇzd´y warp obsahuje stejn´y poˇcet vl´aken8, kter´a jsou pak prov´adˇena multiprocesorem. Aktivn´ım warp˚um jsou v z´avislosti na ˇcase pl´anovaˇcem pˇrep´ın´any pro maximalizaci v´ypoˇcetn´ıho v´ykonu. [5]
• Blok vl´aken (Block)
Blok je skupina warp˚u, kter´a je vykon´av´ana na jednom multiprocesoru. Jeden kernel m˚uˇze b´yt prov´adˇen nˇekolika bloky vl´aken, proto jsou tyto bloky jeˇstˇe sdruˇzeny v tzv.
gridu. Stejnˇe jako u vl´aken, kaˇzd´y blok m´a pˇriˇrazenu jednoznaˇcnou identifikaci v r´amci gridu. Tato identifikace je pˇr´ıstupn´a ve vestavˇen´e promˇenn´e blockIdx a rovnˇeˇz lze z´ıskat i ”rozmˇery” dan´eho bloku, ty jsou obsaˇzeny ve vestavˇen´e promˇenn´eblockDim. U blok˚u je poˇzadov´ano, aby mohly b´yt prov´adˇeny v n´ahodn´em poˇrad´ı, a tedy paralelnˇe.
• Grid
Grid9 je tedy seskupen´ı blok˚u v r´amci jednoho kernelu. Grid je obvykle jedno ˇci dvourozmˇern´y, jeho rozmˇery jsou specifikov´any v prvn´ım parametru v z´avork´ach
<<<...>>> pˇri vol´an´ı kernelu. N´asleduje pˇr´ıklad zdrojov´eho k´odu CUDA pro souˇcet dvourozmˇern´ych matic s vyuˇzit´ım dˇr´ıve popsan´ych struktur.
//Definice kernelu
__global__ void MatAdd(float A[N][N], float B[N][N],float C[N][N]){
int i = blockIdx.x * blockDim.x + threadIdx.x;
int j = blockIdx.y * blockDim.y + threadIdx.y;
if (i < N && j < N)
dim3 dimGrid((N + dimBlock.x - 1) / dimBlock.x, (N + dimBlock.y - 1) / dimBlock.y);
MatAdd<<<dimGrid, dimBlock>>>(A, B, C);
}
Pr´avˇe popsan´a hierarchie vl´aken je uvedena na obr´azku 4.5. Nyn´ı se zamˇeˇr´ım na pamˇet’ov´y model multiprocesoru GPU.
• Registry
Registry pˇredstavuj´ı nejrychlejˇs´ı pamˇet’ na jednom multiprocesoru. Tento typ pamˇet´ı je pˇr´ıstupn´y pouze pro dan´e vl´akno. S t´ım souvis´ı fakt, ˇze ˇzivotnost dat je pˇr´ımo nav´az´ana na dobu ˇzivota dan´eho vl´akna. V uv´adˇen´e hierarchii jsou tedy registry GPU tou nejrychlejˇs´ı pamˇet´ı, avˇsak maj´ı t´eˇz nejmenˇs´ı kapacitu10
8warp-size
9Nab´ız´ı se i ˇcesk´e oznaˇcen´ı - mˇr´ıˇzka
10Stejnˇe je tomu v pˇr´ıpadˇe CPU.
• Sd´ılen´a pamˇet’
Sd´ılen´a pamˇet’ je pamˇet’ pˇr´ıtomn´a na ˇcipu, proto je rychlost sd´ılen´e pamˇeti rovnˇeˇz velmi vysok´a. Do sd´ılen´e pamˇeti maj´ı pˇr´ıstup vˇsechny procesory bloku jak pro z´apis, tak pro ˇcten´ı. Dokumentace uv´ad´ı, ˇze pro vˇsechna vl´akna v jednom warpu je pˇr´ıstup do sd´ılen´e pamˇeti t´emˇeˇr stejnˇe rychl´y jako pˇr´ıstup do registr˚u. Nesm´ı ovˇsem nastat konflikty mezi jednotliv´ymi pamˇet’ov´ymi banky, coˇz jsou ˇc´asti sd´ılen´e pamˇeti, kter´e maj´ı stejnou velikost a jsou souˇcasnˇe pˇr´ıstupn´e pro vˇsechna vl´akna warpu.
• Pamˇet’ textur a konstant
Obˇe tyto pamˇeti se vyznaˇcuj´ı t´ım, ˇze obˇe slouˇz´ı jako rychl´a vyrovn´avac´ı pamˇet’ mezi procesory a maj´ı n´ızkou pˇr´ıstupovou dobu a velmi dobr´y hit-ratio.
• Glob´aln´ı pamˇet’
Glob´aln´ı pamˇet’ je nejpomalejˇs´ı pamˇet´ı v uv´adˇen´e hierarchii. Je pˇr´ıstupn´a jak pro hostitele (host), tak pro zaˇr´ızen´ı (device). Jej´ı ˇzivotnost je stejn´a jako ˇzivotnost dan´e aplikace. Zdroj [6] uv´ad´ı, ˇze glob´aln´ı pamˇet’ m˚uˇze b´yt aˇz 150x pomalejˇs´ı neˇz pamˇet’
sd´ılen´a a registry.
CUDA rozˇsiˇruje jazyk C o moˇznosti prov´adˇen´ı v´ypoˇct˚u paralelnˇe. K tomu je zapotˇreb´ı de-finovat jist´y druh oznaˇcen´ı funkc´ı a promˇenn´ych. Funkce se mohou prov´adˇet jak na zaˇr´ızen´ı (GPU), tak v hostiteli. Tyto skuteˇcnosti je potˇreba rozliˇsovat, proto uv´ad´ım nejd˚uleˇzitˇejˇs´ı specifika.
• __device__ funkce
Tento kvalifik´ator oznaˇcuje funkci, kter´a m˚uˇze b´yt provedena pouze na dan´em zaˇr´ızen´ı a nelze ji volat z hostitelsk´eho programu, jedin´a moˇznost vol´an´ı je z programu zaˇr´ızen´ı.
Pro tento typ funkc´ı existuj´ı urˇcit´a omezen´ı: nen´ı podporov´ana rekurze, funkce mus´ı m´ıt konstantn´ı poˇcet parametr˚u a nen´ı moˇzn´e, aby tento typ funkc´ı ve sv´em tˇele de-finoval statickou promˇennou. Kvalifik´ator nesm´ı b´yt pouˇzit spoleˇcnˇe s kvalifik´atorem __global__.
• __global__ funkce
T´ımto kvalifik´atorem je oznaˇcena funkce - kernel. Takovouto funkci je moˇzn´e zavolat pouze z hostitelsk´eho programu a je vykon´av´ana v´yhradnˇe v zaˇr´ızen´ı. Pro kernely plat´ı stejn´a omezen´ı jako pro funkce s kvalifik´atorem__device__, tj. omezen´ı t´ykaj´ıc´ı se rekurze, konstantn´ıho poˇctu parametr˚u a vytv´aˇren´ı statick´ych promˇenn´ych. D´ale nen´ı moˇzn´e spoleˇcnˇe pouˇz´ıt kvalifik´ator__host__. Funkce oznaˇcen´e jako__global__ -kernely -mus´ım´ıt n´avratovou hodnotu typuvoid. Vol´an´ı kernelu je asynchronn´ı, tzn.
ˇze vrac´ı ˇr´ızen´ı hostitelsk´emu programu jeˇstˇe pˇred t´ım, neˇz bylo dokonˇceno prov´adˇen´ı dan´eho kernelu.
• __host__ funkce
Tento kvalifik´ator oznaˇcuje funkci, kter´a je vol´ana jen a pouze z hostitelsk´eho pro-gramu a rovnˇeˇz je na hostitelsk´em zaˇr´ızen´ı prov´adˇena. Zaj´ımavost´ı je, ˇze kvalifik´atory __host__a__device__mohou b´yt pouˇzity spoleˇcnˇe v deklaraci funkce, takov´a kon-strukce pak znamen´a, ˇze k´od funkce bude pˇreloˇzen jak pro hostitelskou architekturu, tak pro prov´adˇen´ı v zaˇr´ızen´ı.
• __device__ promˇenn´a
Je vytvoˇrena v pamˇet´ı dan´eho zaˇr´ızen´ı a existuje stejnˇe dlouho jako program. Je pˇr´ıstupn´a ze vˇsech vl´aken gridu a tak´e z hostitele skrze knihovn´ı vol´an´ı.
• __constant__ promˇenn´a
Kvalifik´ator m˚uˇze b´yt pouˇzit spoleˇcnˇe s kvalifik´atorem__device__, deklaruje promˇennou, a to takovou, kter´a je uloˇzena v prostoru pamˇeti konstant. Jej´ı ˇzivotnost je stejn´a jako ˇzivotnost aplikace a je pˇr´ıstupn´a ze vˇsech vl´aken gridu a pˇres knihovn´ı vol´an´ı tak´e z hostitele. Do takov´eto promˇenn´e m˚uˇze pˇriˇrazovat pouze hostitel, opˇet skrze knihovn´ı vol´an´ı.
• __shared__ promˇenn´a
Kvalifik´ator shared m˚uˇze b´yt pouˇzit spoleˇcnˇe s kvalifik´atorem device, deklaruje promˇennou, kter´a je uloˇzena v pamˇeti bloku vl´aken, jej´ı ˇzivotnost je tedy stejn´a jako ˇzivotnost bloku a je pˇr´ıstupn´a pouze a jen z dan´eho bloku vl´aken.
Psan´ı programu s pomoc´ı CUDA se tedy velmi podob´a psan´ı bˇeˇzn´eho programu v jazyce C pouze s uveden´ymi rozˇs´ıˇren´ım. Program napsan´y v CUDA m˚uˇze bˇeˇzet na gra-fick´e kartˇe, nebo lze pouˇz´ıt emulovan´y reˇzim, napˇr. pro ladˇen´ı, nebo v pˇr´ıpadˇe, ˇze nen´ı k dispozici odpov´ıdaj´ıc´ı hardware. Je jasn´e, ˇze pˇri pouˇzit´ı emulovan´eho reˇzimu se dosahuje podstatnˇe menˇs´ıho v´ypoˇcetn´ıho v´ykonu, neˇz kdyby v´ypoˇcet prob´ıhal v grafick´e kartˇe. Proto s´am v´yrobce nedoporuˇcuje emulovan´y reˇzim pouˇz´ıvat hlavnˇe pro v´ykonnostn´ı ladˇen´ı dan´e aplikace.
Kapitola 5
N´ avrh programu v CUDA
Souˇc´ast´ı t´eto diplomov´e pr´ace je vytvoˇren´ı programu s poˇzit´ım prostˇred´ı CUDA, kter´y bude simulovat dˇeje prob´ıhaj´ıc´ı v homogenn´ım a nehomogenn´ım veden´ı. Ve sv´ych experimentech s prostˇred´ım CUDA jsem vych´azel z pˇr´ıklad˚u jednoduch´ych program˚u dodan´ych pˇr´ımo s prostˇred´ım CUDA.
5.1 Poˇ c´ ateˇ cn´ı anal´ yza
Diferenci´aln´ı rovnice popisuj´ıc´ı jednotliv´e elementy veden´ı lze simulovat vyuˇzit´ım spojit´e simulace. Proto bude potˇreba tento algoritmus implementovat. Pseudok´od algoritmu spojit´e simulace ukazuje n´asleduj´ı blok.
inicializace, nastaven´ı poˇc´ateˇcn´ıch podm´ınek integr´ator˚um while(t < tEND){
vyhodnocen´ı vstup˚u integr´ator˚u;
proveden´ı v´ypoˇctu;
t = t+h - posun ˇcasu o dan´y krok h }
Z povahy ˇreˇsen´eho probl´emu vypl´yv´a, ˇze pr´avˇe posloupnost pseudo-pˇr´ıkaz˚u z blokuwhile se bude prov´adˇet paralelnˇe. Uvaˇzujme nyn´ı z´apis programu v TSKL z pˇr´ılohy B. Pokud se povede naj´ıt zp˚usob, aby bylo moˇzn´e jednotliv´e rovnice prov´adˇet nez´avisle na sobˇe, pak bude probl´em paralelizovateln´y.
Proto uvaˇzujme jeden nˇekter´y krok v´ypoˇctu. V z´apisu programu se nevyskytuj´ı rychl´e smyˇcky, kter´e by v´ypoˇcet zkomplikovaly. Nav´ıc, v dobˇe, kdy se budou poˇc´ıtat hodnoty pro dan´y krok v´ypoˇctu jsou k dispozici hodnoty pro krok pˇredchoz´ı1. Proto nen´ı probl´em tento v´ypoˇcet paralelizovat.
V n´avaznosti na prostˇred´ı CUDA bude d˚uleˇzit´ym parametrem poˇcet vl´aken, kter´y v tomto pˇr´ıpadˇe bude reprezentovat poˇcet dvojbran˚u simulovan´eho veden´ı. Je tˇreba si uvˇedomit, ˇze tento postup je moˇzn´y pouze pro simulaci homogenn´ıho veden´ı.
V pˇr´ıpadˇe, ˇze se bude simulovat chov´an´ı veden´ı nehomogenn´ıho, je potˇreba vyˇreˇsit probl´emy nehomogenit, jako napˇr. mˇen´ıc´ı se konstanty.
1Resp. poˇc´ateˇcn´ı hodnoty pro krok prvn´ı v ˇcaset= 0
Obr´azek 5.1: N´avrh datov´ych struktur