Skoči na vsebino

Programski vmesnik CUDA1

Do zdaj smo se naučili, kako je zgrajena grafična procesna enota (GPE), kaj so računske enote (SM) in procesni elementi (jedra CUDA). Spoznali smo kako GPE izvaja programske niti, kako so niti organizirane ter kako se ukazi razporejajo na izvršilne enote. Poleg tega zdaj razumemo hierarhijo pomnilnika GPE. Sedaj je čas, da si ogledamo programski vmesnik CUDA, katerega začetki segajo v leto 2007. Do takrat so bile grafične procesne enote namenjene predvsem obdelavi in izrisevanju 3D scen. Uporaba GPE na drugih problemskih domenah je bila zelo otežena. Programski vmesnik CUDA pa je programerjem omogočil, da računske zmogljivosti GPE izkoristijo za reševanje splošnih računskih problemov, predvsem takih, kjer je prisotnega veliko podatkovnega paralelizma.

Heterogen sistem CUDA

CUDA je platforma za vzporedno računanje in programski vmesnik s skupkom razširitev za jezik C. Programski vmesnik CUDA zagotavlja dve ključni funkcionalnosti pri programiranju grafičnih procesnih enot:

  • Zagon dela kode na GPE ter organizacijo niti na GPE.
  • Prenos podatkov med CPE in GPE ter dostop do pomnilnika na GPE.

Programski vmesnik CUDA omogoča izvajanje aplikacij na heterogenih računalniških sistemih preko preprostega označevanja kode s pomočjo razširitev za jezik C. Heterogeno okolje sestavljajo CPE in GPE, vsak s svojim pomnilnikom povezanim z vodilom PCIe. V heterogenem sistemu imamo enega gostitelja, povezanega z enim ali več pospeševalnimi napravami GPE, vsako s svojim pomnilnikom. Tipičen heterogen sistem je prikazan na spodnji sliki.

CUDA Heterogeneous System

Ločiti moramo med sledečima pojmoma:

  • Gostitelj (angl. Host) — predstavlja CPE in pripadajoč pomnilnik (angl. Host memory),
  • Naprava (angl. Device) — predstavlja GPE in pripadajoč pomnilnik (angl. Device memory).

Program CUDA sestoji iz gostiteljskega programa (angl. host program), ki se izvaja na gostitelju (običajno namizni računalnik s CPE), in enega ali več delov programa, ki jim rečemo ščepec (angl. kernel), ki se izvajajo na napravah GPE.

Izvajalni model CUDA

Ključna sestavina izvajalnega modela CUDA je ščepec - koda, ki se izvaja na napravi GPE. Kot programerji ščepec napišemo kot zaporedni program. Ko se ščepec zažene, se izvajanje samodejno preseli na napravo, kjer se ustvari veliko število niti in vsaka nit izvaja ukaze določene v ščepcu. Organizacija niti je ključen del programiranja v CUDA. Organizacijo niti izvedemo hierarhično v treh ravneh, posamezna nit (angl. Thread), blok niti (angl. Block) in mreža blokov (angl. Grid), kot je prikazano na spodnji sliki.

CUDA grid of thread blocks

Z drugimi besedami, vsak ščepec se izvaja znotraj mreže niti. Ta mreža je razdeljena na bloke, vsak blok pa je dodatno razdeljen na niti. Vse niti izvajajo isti ščepec. To pomeni, da se ščepec izvaja v ogromnem številu primerkov ali instanc - vsaka nit izvaja kodo, ki jo vsebuje ščepec, vendar izvaja operacije na svojih lastnih podatkih.

Ključni koncepti, ki si jih moramo zapomniti so:

  • Mreža (angl. Grid) — Vse niti, ki jih sproži zagon ščepca, se imenujejo mreža. Mreža je sestavljena iz blokov niti. Vse niti v mreži si delijo isti globalni pomnilniški prostor. Velikost mreže je določena s številom blokov. Na primer, mreža velikosti 16 vsebuje 16 blokov niti. Mrežo lahko organiziramo v več dimenzijah (do 3). S tem samo spremenimo indeksiranje posameznih blokov. Če je mreža 1D, so vsi bloki v eni dimenziji in imajo po en indeks (0-15). Če je mreža 2D, je 16 blokov razporejenih v dveh dimenzijah (npr. 4x4) - idetificiramo jih preko dvojnega indeksa ((0,0) - (3,3)).

  • Blok niti (angl. Thread Block) — Mreža je sestavljena iz več blokov niti. Blok niti so vse niti, ki tečejo na isti računski enoti in lahko medsebojno sodelujejo z uporabo

  • sinhronizacije znotraj bloka in

  • skupnega pomnilnika bloka.

Niti iz različnih blokov ne morejo sodelovati med seboj (ta trditev ni čisto resnična, novejše GPE do neke mere omogočajo sodelovanje niti iz različnih blokov).

  • Nit (angl. Thread) — Posamezen instanca ščepca na napravi GPE. Na gostitelju določimo, kako naj se algoritem preslika na mrežo niti in blokov glede na podatke, ki jih bomo obdelovali in zmogljivosti naprave GPE. Programer na preprost način (s pisanjem zaporedne kode) napiše algoritma in se ne obremenjuje z ustvarjanjem ter upravljanjem tisočih niti na GPE. Na primer, da želite na GPE opraviti procesiranje slike, kjer je potrebno izvesti neko operacijo nad vsemi slikovnimi točkami. V tem primeru bo vaš ščepec vseboval kodo za delo s posamezno slikovno točko. Nato boste zagnali 2D mrežo niti, kjer dimenzije mreže ustrezajo velikosti slike. Vsaka nit iz mreže bo obdelala svojo slikovno točko, končen rezultat pa bo v celoti sprocesirana slika.

Izvajalni model CUDA je v osnovi sočasen. Ko zaženemo ščepec, se nadzor takoj vrne gostitelju, s čimer se CPE sprosti za izvajanje nadaljnih nalog. Prav tako se lahko zvajanje na GPE prekriva s komunikacijo med gostiteljem in GPE.

Indeksiranje niti

Identifikacija posameznih niti poteka preko dveh unikatnih koordinat:

  • blockIdx (indeks bloka v mreži) in
  • threadIdx (indeks niti v bloku).

Te spremenljivke so vgrajene in vnaprej inicializirane. Do njih je mogoče dostopati znotraj ščepca. Ko se ščepec zažene, CUDA dodeli koordinate blockIdx in threadIdx vsaki niti. Na podlagi teh koordinat lahko razdelimo dele podatkov različnim nitim. Koordinatni spremenljivki blockIdx in threadIdx sta tipa dim3. Gre za strukturo, ki vsebuje tri nenegativna cela števila. Ta tri števila so dostopna prek polj x, y in z, npr. blockIdx.x, blockIdx.y in blockIdx.z ter threadIdx.x, threadIdx.y in threadIdx.z.

Dimenzije mreže in bloka določata naslednji vgrajeni spremenljivki:

  • blockDim (dimenzija bloka, merjena v nitih) in
  • gridDim (dimenzija mreže, merjena v blokih).

Te spremenljivke so prav tako tipa dim3. Vsaka komponenta v spremenljivki tipa dim3 je dostopna prek polj x, y in z, npr. blockDim.x, blockDim.y in blockDim.z. Pri definiranju spremenljivke tipa dim3 se vsaka nepodana komponenta samodejno inicializira na 1.

CUDA grid

Slika zgoraj prikazuje primer hierarhije in indeksiranje niti razporejenih v 2D mrežo, ki vsebuje 2D bloke. Organizacijo niti je določena s pomočjo zgoraj naštetih spremenljivk, ki so dostopne znotraj ščepca. Organizacija je običajno prilagojena problemu, ki ga rešujemo in je izberemo tako, da si čim bolj olajšamo programiranje ščepcev ter čim bolj izkoristimo GPE. Na sliki imamo postavljeno 2D mrežo blokov velikosti 3x3. Vsak blok je prav tako organiziran v dveh dimenzijah in vsebuje 3x3=9 niti. Vsega skupaj imamo torej 81 niti. S tako organizacijo niti, bi lahko zelo enostavno napisali ščepec, ki sočasno obdela vse slikovne točke neke slike velikosti 9x9. Določiti moramo seveda, katera nit bo obdelala katero slikovno točko. Opraviti moramo preslikavo med hierarhičnim indeksiranjem niti in 2D indeksiranjem slikovnih točk. Kot primer izračunajmo, katero slikovno točko naj obdela nit (1,2), ki se nahaja v bloku (1,0). To storimo na naslednji način: 

x = blockIdx.x x blockDim.x + threadIdx.x = 1 x 3 + 1 = 4
y = blockIdx.y x blockDim.y + threadIdx.y = 0 x 3 + 2 = 2
Izbrana nit bo torej obdelala slikovno točko v stolpcu z indeksom 4 in vrstici z indeksom 2. Če zgornji izračun opravi vsaka nit, na enostaven način razdelimo delo med vse niti. Pred tem moramo seveda poskrbeti, da smo zagnali dovolj veliko mrežo niti, da pokrijemo z njo celotno sliko. Kako to storimo si bomo ogledali v nadaljevanju.

Ščepci in organizacija niti

S pojmom ščepec (angl. kernel) imenujemo del programske kode (funkcijo), ki se bo izvajala na GPE. Ščepec definiramo z uporabo ključne besede __global__ pred deklaracijo funkcije. Ta ključna beseda označuje, da funkcijo kličemo iz gostitelja, vendar se bo izvedla na napravi (GPE). Ščepec vedno vrača podatkovni tip void. Primer ščepca:

1
2
3
4
5
// kernel function declaration
__global__ void kernelFunc (type1 parameter1, type2 parameter2, ...)
{
 // function body implements the task for one thread
}

Pri klicu ščepca moramo povedati tudi organizacijo niti. Navesti moramo število blokov, ki jih želimo ter število niti v bloku:

1
2
// launch the device computation kernel
kernelFunc <<< grid_size , block_size >>>( parameter1 , parameter2, ... );
Ščepec se vedno začne izvajati sočasno s kodo, ki teče na gostitelju. To pomeni, da lahko CPE nadaljuje z drugimi opravili, medtem ko čaka na rezultate iz GPE. Pri določanju velikosti mreže lahko posamezne dimenzije izpustimo. Za zgornji primer obdelave slike velikosti 9x9 slikovnih točk bi mrežo niti organizirali na naslednji način:

1
2
dim3 grid_size(3, 3);
dim3 block_size(3, 3);

Poleg ščepcev, ki so funkcije, klicane iz gostitelja in izvajane na napravi, lahko v CUDA pišemo tudi druge vrste funkcij. Možni kvalifikatorji vrst funkcij, so:

  • __global__: funkcije se izvajajo na napravi in jih kličemo iz gostitelja. Za naprave s Compute Capability >= 3.5 lahko takšne funkcije kličemo tudi iz naprave.
  • __device__: funkcije se izvajajo na napravi in jih kličemo samo iz naprave.
  • __host__: funkcije se izvajajo na gostitelju in jih kličemo iz gostitelja, ta kvalifikator lahko tudi izpustimo, saj je to privzeto obnašanje.

Pri funkciji lahko uporabimo kvalifikatorja device in host skupaj. V tem primeru je funkcija prevedena tako za gostitelja kot za napravo.

Pomnilniški model CUDA

Programski model CUDA predvideva sistem, ki je sestavljen iz gostitelja in naprave, vsak s svojim ločenim pomnilnikom: gostiteljskim pomnilnikom in pomnilnikom naprave. CUDA nudi funkcije za rezervacijo pomnilnika naprave, sproščanje pomnilnika naprave ter prenos podatkov med gostiteljem in pomnilnikom naprave. Poleg tega je skozi programski model razkrita hierarhija pomnilnika GPE, ki je prikazana na spodnji sliki.

CUDA abstraction of memory hierarchy

Na dnu slike vidimo globalni pomnilnik in konstantni pomnilnik. To sta pomnilnika, v katera gostitelj (CPE) lahko zapisuje in bere. Vse niti v mreži lahko dostopajo do globalnega pomnilnika, vendar ima globalni pomnilnik največjo zakasnitev pri dostopu. Konstantni pomnilnik omogoča samo branje s strani naprave. Fizično je skupaj z globalnim pomnilnikom, vendar je predpomnjen ločeno in zato običajno hitrejši kot slednji. Vsaka nit ima svoj nabor registrov, niti iz bloka pa si delijo skupni pomnilnik, ki je lasten vsakemu bloku niti.

Do registrov in skupnega pomnilnika je mogoče dostopati z zelo visoko hitrostjo na izjemno vzporeden način. Vendar pa majhna velikost teh pomnilnikov predstavlja oviro pri obdelavi večjih količin podatkov. Ščepec običajno uporablja registre za hranjenje pogosto dostopanih spremenljivk, ki so zasebne za vsako nit. Vse niti v bloku lahko dostopajo do spremenljivk hranjenih v skupnem pomnilniku. Z njegovo pomočjo je možno vzpostaviti sodelovanje med nitmi istega bloka.

Upravljanje z globalnim pomnilnikom

Rezervacijo pomnilnika na GPE izvedemo s pomočjo klica funkcije cudaMalloc s podpisom:

cudaError_t cudaMalloc(void** devPtr, size_t count)

Funkcija rezervira count bajtov globalnega pomnilnika na napravi ter vrne lokacijo tega pomnilnika v kazalcu devPtr. Nikoli se ne smemo poskušati sklicevati na ta pomnilnik iz gostitelja preko tega kazalca! Pomnilniška prostora gostitelja in naprave sta ločena. Tako rezerviran pomnilniški prostor ni inicializiran. Programer je odgovoren, da dodeljeni globalni pomnilnik napolne s podatki prenešenimi iz gostitelja.

Ko aplikacija ne potrebuje več dodeljenega globalnega pomnilnika, ga lahko sprosti z uporabo funkcije:

cudaError_t cudaFree(void *devPtr)

Ta funkcija sprosti globalni pomnilnik, na katerega kaže devPtr. Dodeljevanje in sproščanje pomnilnika na napravi sta dragi operaciji, zato vedno skušamo minimizirati število klicev teh funkcij.

Ko je globalni pomnilnik dodeljen lahko podatke prenesemo na napravo iz gostitelja. Funkcija, ki se uporablja za prenos podatkov med gostiteljem in napravo je:

cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind)

Ta funkcija kopira count bajtov iz lokacije pomnilnika src na lokacijo pomnilnika dst, v smeri, določeni s kind, kjer kind zavzame eno od naslednjih vrednosti:

  • cudaMemcpyHostToHost,
  • cudaMemcpyHostToDevice,
  • cudaMemcpyDeviceToHost,
  • cudaMemcpyDeviceToDevice.

Funkcija je blokirajoča, saj počaka, da se prenos v celoti zaključi, šele nato se lahko izvajanje na CPE nadaljuje.

Upravljanje s skupnim pomnilnikom

Ščepec lahko eksplicitno upravlja s vsebino skupnega pomnilnika, zato ga pogosto opisujejo kot programsko upravljani predpomnilnik. Programer ima popoln nadzor nad tem, kdaj se podatki premaknejo v skupni pomnilnik in kdaj se podatki izrinejo.

Pri dekalraciji pomnilniške strukture za katero želimo, da se hrani v skupnem pomnilniku uporabimo kvalifikator shared. Naslednji odsek kode statično deklarira 2D celoštevilsko tabelo matrix in jo postavi v skupni pomnilnik.

__shared__ int matrix[10][10];

Ta način ima očitno omejitev, saj moramo velikost podatkovne strukture poznati vnaprej. Možna je tudi dinamična rezervacija skupnega pomnilnika pri zagonu ščepca. To storimo tako, da pri zagonu poleg organizacije niti dodamo še en parameter, ki določa velikost podatkovne strukture, ki jo bomo hranili v skupnem pomnilniku v bajtih. Primer: 

__global__ void kernelFunc (type1 parameter1, type2 parameter2, ...)
{
    __shared__ int array[];
}
//...
kernelFunc<<<grid_size, block_size, shared_mem_size>>>(argumet1, argument2, ...);
//...

  1. © Patricio Bulić, Davor Sluga Univerza v Ljubljani, Fakulteta za računalništvo in informatiko. Gradivo je objavljeno pod licenco Creative Commons Priznanje avtorstva-Nekomercialno-Deljenje pod enakimi pogoji 4.0 Mednarodna. ↩