OpenCL – Deo I
GPU računarstvo predstavlja koršćenje grafičke procesorske jedinice za izračunavanja opšte namene umesto tradicionalnog grafičkog prevođenja. GPU je višejezgarni multiprocesor visokih performansi koji služi da ubrza razne aplikacije koristeći paralelno računarstvo.
Sadržaj
Kratka istorija GPU računarstva
GPU kompjuterska revolucije je započeta u novembru 2006. godine, kada je AMD uveo Close to Metal (CTM), nizak sloj hardverskog i programerskog interfejsa koji dozvoljava programerima da iskoriste prednosti skupa domaćih instrukcija i memorije modernog GPU-a za izračunavanja opšte namene. CTM omogućava programerima da kreiraju kompajlere, dibagere, biblioteke i aplikacijske platforme. Uvodjenjem CTM-a, kreirana je nova klasa aplikacija.
Kako bi se što bolje iskoristili AMD GPU procesori, 2007. godine AMD je kreirao ATI Stream SDK v1. ATI Stream SDK v1 je dodao jezik visokog nivoa poznat kao ATI Brook+. CTM se razvio u ATI CAL (Compute Abstraction Layer – Kompjuterski apstraktni sloj), podržavajući API sloj za Brook+. Uvođenjem ATI Stream SDK v1, AMD je omogućio alat niskog i visokog nivoa za pristup opšte namene AMD GPU hardveru.
Mana korišćenja ATI Stream SDK v1 je da aplikacije napravljene s SDK-om mogu da se pokrenu samo na AMD GPU hardveru.
U junu, 2008. godine, AMD je s drugim industrijskim igračima u GPU računarstvu osnovao OpenCL radnu grupu od strane Kronus grupe (The Khronos Group). Khronos, poznat kao vođa u kreiranju drugih otvorenih specifikacija kao što je OpenGL, je bio izabran da kreira i specifikaciju za OpenCL. Posle 5 meseci završena je i publikovana tehnička dokumentacija za OpenCL 1.0. AMD je najavio da će usvojiti OpenCL programski standard i da će izvršiti integraciju kompatibilnog kompajlera s izvršavanjem u besplatan ATI Stream SDK v2. U decembru 2009, AMD je imao ATI Stream SDK v2 s podrškom OpenCL-a 1.0.
Heterogeno računarstvo
Heterogeno računarstvo uključuje razne vrste računarskih jedinica. Računarska jedinica može biti jedinica obrade opšte namene, kao što je centralna procesorska jedinica (CPU), grafička procesorska jedinica (GPU), ili jedinica obrade specijalne namene kao što je procesor digitalnih signala (DSP).
Modernim kompjuterski aplikacijama kojima je potrebna interakcija s raznim sistemima (kao što su audio/video sistemi, umrežene aplikacije, i drugo) napredak u CPU tehnologiji nije bio dovoljan. Bolje performanse je obezbeđivao specijalizovan hardver, koji je činio sistem heterogenim. Dodavanje raznih vrsta računarskih jedinica u heterogene sisteme omogućavalo je dizajnerima aplikacija da izaberu najpogodniju jedinicu na kojoj će se zadatak izvršavati.
Šta je OpenCL?
OpenCL (The Open Computing Language) je besplatan i paralelan računarski API dizajniran da omogući GPU i drugim koprocesorima da rade zajedno s CPU-om. Kao standard, OpenCL 1.0 je kreiran 8. decembra 2008. godine, od strane Kronus grupe. Programeri su dugo čekali na mogućnost da podele problem na više konkurentnih podskupova, čineći ih izvodljivim na GPU-u koji je iskorišćen kao matematički koprocesor zajedno s CPU-om.
Primarna dobit od OpenCL-a je znatno ubrzanje paralelne obrade. OpenCL koristi računarske resurse kao što su višejezgarni CPU-i i GPU-i. OpenCL takođe dopunjuje postojeći OpenGL vizuelni API deljenjem struktura podataka i memorijskih lokacija.
Ovaj sloj niskog nivoa pravi granicu između hardvera i gornjeg sloja softvera. Sve specifičnosti hardverske implementacije, kao što su driveri i izvršavanje, nisu vidljive programerima gornjih softverskih slojeva zbog korišćenje abstrakcija visokog nivoa. Zbog toga programeri koriste najbolji hardver bez uplitanja infrastukture gornjeg softvera. Promena od vlasničkog standarda ka otvorenom standardu takođe doprinosi ubrzanju opštih izračunavanja.
OpenCL razvojni okvir je sastavljen od tri glavna dela:
- Specifikacije jezika
- Platformskog sloja API-ja
- API-ja izvršavanja
Specifikacija jezika opisuje sintaksu i progamski interfejs za pisanje jezgarnih (kernel) programa koji se pokreću na podržanim jedinicama, kao što su GPU, višejezgarni CPU, ili DSP. Jezgra (kerneli) mogu biti prekompajlirana ili OpenCL izvršavanje kompajlira jezgro programa tokom izvršavanja.
OpenCL programerski jezik se bazira na ISO C99 specifikaciji s dodatim proširenjima i ograničenjima. Dodaci uključuju vektorske tipove i vektorske operacije, optimizovan pristup slikama, kvalifikator adresnog prostora. Ograničenja ne uključuju podršku za pokazivače funkcija, bit polja, i rekurzije. Prva baza OpenCL-a je bio C jezik zbog svoje rasprostranjenosti.
Platformski sloj API-ja omogućava programerima pristup rutinama softverskih aplikacija koji mogu da traže od sistema OpenCL-podržane uređaje. Ovaj sloj omogućava programerima korišćenje koncepta konteksta uredjaja i radnih redova. Omogućava im da izaberu i inicijalizaciju OpenCL uređaje, da pošalju posao uređajima i izvrše prenos podataka od i do uređaja.
API izvršavanje koristi kontekst za upravljanje objektima kao što su komandni redovi, memorijski objekti i objekti jezgra. Podjednako dobro koristi kontekst za izvršavanje jezgra na jednom ili više uređaja određenih kontekstom.
Platformski model
OpenCL platformski model se definiše kao host povezan s jednim ili više OpenCL uređaja. Slika 2 prikazuje platformski model koji obuhvata jedan host i više računarskih uređaja, od kojih svaki sadrži više procesorskih elemenata.
Host je bilo koji kompjuter s pokrenutim CPU-om na kojem radi standardni operativni sistem. Uređaji OpenCL-a mogu biti GPU, DSP, ili višejezgrani CPU. Svaki OpenCL uređaj sadrži jednu ili više kompjuterskih jedinica (jezgara). Kompjuterska jedinica sadrži jednu ili više procesorskih elemenata. Procesorski elementi izvršavaju instrukcije kao SIMD (Single Instruction, Multiple Data – jedna instrukcija, više podataka) ili SPMD (Single Program, Multiple Data – jedan program, više podataka). SPMD instrukcije se obično izvršavaju na uređajima opšte namene kao što je CPU, dok SIMD instrukcije zahtevaju vektorski procesor kao što je GPU.
Sledeći prikaz (Slika 3) ATI Radeon HD 5870 GPU arhitekture prikazuje konstrukciju komputerskog uređaja. ATI Radeon HD 5870 GPU je sastavljen od 20 SIMD jedinica, koje prevode 20 kompjuterskih jedinica u OpenCL-u:
Svaka SIMD jedinica se sastoji od 16 jezgara protoka (slika 4). Svako jezgro ima 5 procesorskih elemenata (slika 5). Dakle, svaka kompjuterska jedinica u ATI Radeon HD 5870 GPU-u ima 80 (16 x 5) elemenata obrade.
Model izvršavanja
OpenCL model izvršavanja se sastoji od dve komponente: jezgra i programa host-a. Jezgra čine osnovnu jedinicu izvršnog koda koji se izvršava na jednom ili više OpenCL uređaja. Jezgra su slična C funkcijama čiji podaci ili zadaci mogu biti paralelizovani. Host program se izvršava na host sistemu. Host program definiše kontekst uređaja i ređa kernel izvršne instance koristeći komande redove. Jezgra čekaju u redu, ali mogu biti izvršena po redu ili bez redosleda.
OpenCL koristi paralelno izračunavanje na računarskim uređajima prevodeći problem u N-dimenzionalni indeksni prostor. Kada host program postavi jezgro na čekanje na izvršavanje definiše se indeksni prostor. Svaki nezavisni element izvršavanja u ovom indeksom prostoru se naziva radnom jedinicom (work-item). Svaka radna jedinica izvršava istu funkciju jezgra, ali nad različitim podacima. Kada se jezgro postavi u komandni red, indeksni prostor mora biti definisan da bi uređaj pratio ukupan broj radnih jedinica koje zahtevaju izvršavanje. N-dimenzionalni indeksni prostor može biti \(\) $$ N \in {1,2,3} $$ . Procesiranje lineranog niza podataka ima N=1; procesiranje slika ima N=2; procesiranje 3D zapremine ima N=3.
Procesiranje slike 1024×1024 bi bilo obrađeno na ovaj način: Globalni indeks prostor bi se sastojao od 2-dimenzionalnog prostora 1024×1024. Postoji jedno izvršavanje jezgra (ili radne jedinice) za svaki pixel, tj. 1.048.576 izvršavanja. U ovom indeksnom prostoru, svakoj radnoj jedinici se dodeljuje jedinstven globalni ID. Radna jedinica čija je tačka x=30, y=22 bi mogla da ima globalni ID (30,22).
OpenCL grupiše radne jedinice u radne grupe, kako je prikazano u narednoj slici 6. Veličina svake radne grupe je definisana sopstvenim lokalnim indeks prostorom. Sve radne jedinice unutar jedne grupe se izvršavaju zajedno na istom uređaju. Razlog zbog kojeg se izvršavaju na istom uređaju je dozvola radnim jedinicama da dele lokalnu memoriju i da se sinhronizuju. Globalne radne jedinice su nezavisne i ne mogu se sinhronizovati. Sinhronizacija je jedino dozvoljena između radnih jedinici u radnoj grupi.
Sledeći primer (Slika 7) prikazuje 2-dimenzionalnu sliku s globalnom veličinom od 1024 (32×32). Indeksi prostor je podeljen na 16 radnih grupa. Izdvojena radna grupa ima ID (3,1) i lokalnu veličinu od 64 (8×8). Izdvojena radna jedinica u radnoj grupi ima lokalni ID (4,2), ali može biti adresirana i na osnovu globalnog ID (28,10).
Scalar C function
void square(int n, const float *a, float *result) { int i; for (i = 0; i < n; i++) result[i] = a[i] * a[i]; }Data-Parallel Function
kernel void dp_square (global const float *a, global float *result) { int id = get_global_id(0); result[id] = a[id] * a[id]; } // dp_square execute over "n" work-itemsPrimer iznad prikazuje kako se kernel u OpenCL-u implementira. U ovom primeru svaki element u nizu je kvadriran. Skalarna funkcija bi sadržala prostu for petlju koja bi iterirala kroz sve elemente niza i kvadrirala ih. Pristup u kojem se paralelizuju podaci čita svaki element iz niza paralelno, izvršava operaciju paralelno i ispisuje rezultat na izlaz. U odeljku Data-Parallel Function je kod u kojem se prvo čita indeks vrednost za određenu instancu jezgra, izvršava se operacija i ispisuje se.
Host program je odgovoran za postavljanje i upravljanje izvršavanjem jezgara na OpenCL uređajima korišćenjem konteksta. Koristeći OpenCL API, host kreira i upravlja kontekstom uključujući sledeće resurse:
- Uređaje – Skup OpenCL uređaja koje host koristi da bi izvršio jezgra
- Objekte programa – Programski objekat koji implementira jezgro ili kolekciju jezgara
- Jezgra – Specifične OpenCL funkcije koje se izvršavaju na OpenCL uređajima
- Memorijski objekti – Skup memorijskih bafera ili memorjiskih mapa namenjenih hostu ili OpenCL uređajima.
Kada se kontekst kreira, kreiraju se komandni redovi, koji upravljaju izvršavanjem jezgara na OpenCL uređajima koji su povezani s kontekstom. Komandni redovi prihvataju tri vrste komandi:
- Izvršne komande jezgra pokreću komande jezgra na OpenCL uređajima
- Memorijske komande prenose memorijske objekte između memorijskog prostora hosta i memorijskog prostora OpenCL uređaja
- Komande sinhronizacije definišu redosled u kojem se komande izvršavaju
Komande su smeštene u komandni red i mogu se izvršavati po redu ili bez redosleda. Ako se komande izvršavaju po redu, izvršavaju se u redosledu u kojem su postavljene u red. U suprotnom, redosled izvršavanja se bazira na ograničenjima sinhronizacije.
Memorijski model
Pošto zajednički memorijski prostor nije dostupan host-u i OpenCL uređajima, OpenCL memorijski model definiše četri regiona memorije. Definisani regioni su dostupni radnim jedinicama kada se jezgro izvršava. Slika 8 prikazuje regione memorije dostupne host-u i kompjuterskim uređajima.
Globalna memorija je memorijski region u kojem sve radne jedinice i radne grupe imaju pristup za pisanje i čitanje na kompjuterskom uređaju i na hostu. Ovaj region memorije može biti alociran samo od strane host-a tokom izvršavanja.
Konstantna memorija je region globalne memorije u koji ostaje konstantan tokom izvršavanja jezgra. Radne jedinice imaju samo pristup čitanja u ovom regionu. Hostu je dozvoljen pristup čitanja i pisanja.
Lokalna memorija je memorijski region korišćen za deljene podatke od strane radnih jedinica u radnoj grupi. Sve radne jedinice u istoj radnoj grupi imaju pristup pisanja i čitanja.
Privatna memorija je region dostupan samo za jednu radnu jedinicu.
Memorija hosta i memorija uređaja su nezavisne jedna od druge. Upravljanje memorijom omogućava deljenje podataka između hosta i kompjuterskog uređaja. To znači da podaci moraju biti preneti iz memorije host u globalnu memoriju, pa u lokalnu memoriju i nazad. Ovaj proces radi tako što se komande čitanja/pisanja ređaju u komandni red. Ove komande mogu biti blokirajuće ili neblokirajuće. Blokirajuće komande podrazumevaju da komande host-a čekaju dok se memorijska transakcija izvrši pre nego što nastave s radom. Neblokirajuće komande host postavlja u komandni red i nastavlja s radom, ne čeka da se memorijska transakcija završi.
