OpenCL – Deo II

U prethodnom tekstu je data kratka istorija GPU računarstva, uveden je pojam OpenCL-a i teorijskih modela OpenCL razvojnog okvira. Ovaj članak opisuje strukturu OpenCL programa i veze između uređaja, konteksta, programa, jezgara, memorijskih objekata i komandnih redova.

Izvršavanje OpenCL programa

OpenCL okvir je podeljen na platformski sloj API-ja i API izvršavanja. Platformski API omogućava aplikacijama da traže OpenCL uređaje i da njima upravljaju. API izvršavanja upravlja izvršavanjem jezgra na OpenCL uređajima. Osnovini koraci koji kreiraju program su predstavljeni na slici 1.

Izvršavanje OpenCL programa zahteva:

1. Obezbeđivanje OpenCL uređaja
2. Kreiranje konteksta koji je povezan s OpenCL uređajima
3. Kreiranje programa koji će se pokrenuti na jednom ili više uređaja
4. Biranje jezgra koja će se izvršiti
5. Kreiranje memorijskih objekata na hostu ili na uređaju
6. Kopiranje memorijskih podataka na uređaj kada je potrebno
7. Obezbeđivanje argumenata za jezgra
8. Slanje jezgra u komandni red za izvršavanje
9. Kopiranje rezultata s uređaja na host

Postavljanje resursa

Ova sekcija opisuje kako obezediti OpenCL resurse kao što su informacije o platformi, uređajima, kreiranje konteksta i komandnih redova.

Prvi korak u svakoj OpenCL aplikaciji je zahtevanje informacija o platformi i OpenCL uređajima. Funkcija clGetPlatformIDs() služi da vrati listu dostupnih platformi, kao što je prikazano u sledećem kodu:

cl_platform_id platforme; cl_uint broj_platformi; // traži jednu dostupnu platformu cl_int greska = clGetPlatfromIDs( 1,                    // broj ulaza koji mogu biti dodati na id_platforme &platforme,           // lista pronađenih OpenCL-a &broj_platformi);      // broj pronađenih OpenCL platformi 

clGetPlatformIDs vraća sledeće vrednosti:

• CL_INVALID_VALUE – platforme i broj_platformi su NULL vrednosti
• CL_SUCCESS – Funkcija se uspešno izvršila

Funkcija clGetDeviceIDs() se koristi da pretragu uređaja u sistemu. Sledeći kod traži jedan GPU uređaj:

cl_device_id id_uredjaja; cl_uint broj_uredjaja; cl_int greska; greska = clGetDeviceIDs( id_platforme,          // id_platforme dobijen iz funkcije clGetPlatformIDs CL_DEVICE_TYPE_GPU,   // tip traženog uređaja 1,                    // broj id-a dodatog na id_uredjaja &id_uredjaja,           //lista  id-a uređaja &broj_uredjaja // broj pronađenih uređaja ); 

clGetDeviceIDs() vraća vrednosti:

• CL_INVALID_PLATFORM – Platforma nije ispravna.
• CL_INVALID_DEVICE_TYPE – Uređaj nije ispravna vrednost
• CL_INVALID_VALUE – broj_uredjaja i uređaji su NULL
• CL_DEVICE_NOT_FOUND – Nije pronađen ni jedan odgovarajući uredjaj (device_type)
• CL_SUCCESS – Funkcija je izvršena uspešno

Mogući su sledeći tipovi uređaja (Cl_device_type):

1. CL_DEVICE_TYPE_CPU – OpenCL uređaj je procesor hosta. Host procesor pokreće OpenCL implementaciju i može biti jednojezgaran ili višejezgaran.
2. CL_DEVICE_TYPE_GPU – OpenCL uređaj je GPU.
3. CL_DEVICE_TYPE_DEFAULT – Podrazumevani OpenCL uređaj u sistemu
4. CL_DEVICE_TYPE_ALL – Svi OpenCL uređaji dostupni u sistemu

Kada se identifikuju id_uredjaja, oni mogu da se povežu s kontekstom. OpenCL API izvršavanja koristi kontekst da upravlja komandnim redovima, programskim objektima, objektima jezgra, objektima deljene memorije za uređaje povezane s kontekstom. Funkcija clCreateContext() se koristi za kreiranje konteksta.

cl_context contekst; // lista konteksta, koja mora biti završena nulom karakteristike[0]= CL_CONTEXT_PLATFORM; // definisanje platforme karakteristike[1]= (cl_context_properties) id_platforme; karakteristike[2]= 0; contekst = clCreateContext( karakteristike,      // lista karakteristika konteksta 1,               // broj uređaja u listi id_uredjaja &id_uredjaja,      // lista id_uredjaja NULL,            // pokazivač na grešku callback funkcije NULL,            // argument podataka koji se prosleđuje callback funkciji &greska             // povratni kod ); 

Podržana karakteristika je:

CL_CONTEXT_PLATFORM – njen tip je cl_platform_id. Određuje platformu koja se koristi.

Ako je kontekst uspešno kreiran funkcija clCreateContext() vraća vrednost različitu od nule s greškom čija je vrednost CL_SUCCESS. Inače NULL je vraćeno i jedno od sledećih grešaka:

• CL_INVALID_PLATFORM – Lista karakteristika je NULL ili vrednost platforme nije ispravna

• CL_INVALID_VALUE –

  1. Ime karakteristike u listi karakteristika nije ispravno
  2. Broj uređaja je nula.
  3. Id_uredjaja je NULL.
  4. Uredjaj u listi id_uredjaja nije ispravan ili nije povezan s platformom.

• CL_DEVICE_NOT_AVAILABLE – uredjaj u listi id_uredjaja je nedostupan

• CL_OUT_OF_HOST_MEMORY – Host je onemogućen da alocira OpenCL resurse

Kada je kontekst utvrđen, kreiraju se komandni redovi kojim se omogućava slanje komande uređaju koji je povezan s kontekstom. Funkcija clCreateCommandQueue() kreira red za komande:

cl_command_queue red_za_komande; red_za_komande = clCreateCommandQueue( kontekst,   //ispravan kontekst id_uredjaja, // ispravan uredjaj povezan s kontekstom 0,         // karakteristika za red, koja nije iskorišćena ovde &greska       // povratni kod ); 

Ako kontekst sadrži CPU i GPU uređaje, potrebno je kreirati odvojene redove za komande.

Ako se funkcija uspešno izvrši vraća se vrednost CL_SUCCESS. Inače vraća se NULL uz jednu od propratnih grešaka:

1. CL_INVALID_CONTEXT – Kontekst nije ispravan
2. CL_INVALID_DEVICE – Uređaj nije ispravan ili nije povezan s kontekstom
3. CL_INVALID_VALUE – Lista karakteristika nije ispravna
4. CL_INVALID_QUEUE_PROPERTIES – Uređaj ne podržava karakteristike specificirane u listi karakteristika
5. CL_OUT_OF_HOST_MEMORY – Host ne može da alocira OpenCL resurse.

OpenCL program je sastavljen od funkcija jezgra. Kernel funkcija se identifikuje kvalifikatorom __kernel u programskom kodu.

Objekti programa sadrže programske kodove ili binarne fajlove, poslednji izvršni program, listu uređaja na kojoj se program izvršava.

__kernel void pozdrav_svete(__global char* ulazi, __global char* izlazi) { int broj = get_global_id(0); izlazi[broj] = ulazi[broj] + 1; } 

Za kreiranje objekta programa koristi se funkcija clCreateProgramWithSource():

cl_program program; program = clCreateProgramWithSource( kontekst,           // ispravan kontekst 1,                   //broj stringova u sledećem parametru (const char **) &kodPrograma, // niz stringova NULL,               // dužina svakog stringa ili NULL &greska               // vrednost greške ); 

U ovom primeru je kod programa učitan u niz stringova. Ako je programski kod u odvojenom fajlu, onda mora da se učita fajl u string i da se pošalje funkciji clCreateProgramWithSource(). Moguće je učitati više takvih fajlova i oni bi se smestili u niz stringova. U tom slučaju se mora dodati i dužina niza.

Ako se funkcija uspešno izvrši vraća se vrednost CL_SUCCESS. Inače vraća se NULL uz jednu od propratnih grešaka:

1. CL_INVALID_CONTEXT – Kontekst nije ispravan
2. CL_INVALID_VALUE – Dužina stringa je nula ili niz sadrži NULL vrednost
3. CL_OUT_OF_HOST_MEMORY – Host ne može da alocira OpenCL resurse.

Od ove tačke program je spreman za kompajliranje i linkovanje pre nego što se izvrši na uređaju.

Sledeći korak je kompajliranje i povezivanje programskih objekata funkcijom clBuildProgram():

greska = clBuildProgram( program, // ispravan programski objekat 0,       // broj uređaja liste uređaja NULL,    // lista uređaja – NULL označava sve uređaje NULL,    // string za opciju kreiranja NULL,    // callback funkcija kada se izvršni program iskreira NULL     // argumenti podataka za callback funkciju ); 

Funkcija vraća CL_SUCCESS ako se uspešno izvršila.

Objekat jezgra se kreira posle kreiranja objekta programa. Objekat jezgra je poslat redu za komande za izvršavanje. Funkcija clCreateKernel() kreira objekte jezgra:

cl_kernel jezgro; jezgro = clCreateKernel( program,  // uspešno kreirani objekat programa "hello",     // ime jezgra deklarisanog s __kernel &greska   // kod greške ); 

Ako se funkcija uspešno izvršila vraća objekat jezgra s vrednošću greške CL_SUCCESS.

Pre nego što se objekat jezgra pošalje komandnom redu, treba obezbediti ulazne i izlazne bafere za sve argumente __kernel funkcije.

Argumenti jezgra se postavljaju funkcijom clSetKernelArg().

greska = clSetKernelArg( jezgro,          // ispravan objekat jezgra 0,               // određen argument indeksa jezgra sizeof(cl_mem),  // veličina podatka argumenta &ulazni_podaci      // pokazivač na podatke korišćene kao argumente ); 

Sledeći kod prikazuje deklaraciju __kernel funkcije:

__kernel void hello(__global float *ulazi, __global float *izlazi) 

Izvršavanje programa

U prethodnom tekstu je objašnjeno da OpenCL koristi paralelno izračunavanje na računskim uređajima prevodeći problem u N-dimenzionalni indeksni prostor. Svako izvršavanje jezgra u OpenCL-u se naziva radnom jedinicom. OpenCL koristi paralelno izračunavanje tako što izvršava jezgro na različitim delovima N-dimenzionalnog prostora.Svaka radna jedinica se izvršava s svojim podacima. Važno je odrediti koliko je radnih jedinica potrebno za procesiranje svih podataka. Pre nego što se odredi broj radnih jedinica, treba da se odredi dimenzija N koja predstavlja podatke. Procesiranje lineranog niza podataka ima N=1; procesiranje slika ima N=2; procesiranje 3D zapremine ima N=3.

Na slici 2 je prikazana reprezentacija N-dimenzionalnog prostora.Kada je prostor određen, zbir radnih jedinica (globalna veličina posla) se može preračunati. Analogno, se preračunava lokalna veličina posla (veličina radne grupe). Za jednodimenzionalni skup podataka globalna veličina posla je dužina niza, za dvodimenzionalni skup podataka – kao što je slika – globalna veličina posla je proizvod širine i dužine u pixelima. Za trodimenzionalni skup podataka globalna veličina posla xyz.

OpenCL omogućava radnim jedinicama da se kombinuju u radne grupe. Sve radne jedinice u radnoj grupi se izvršavaju na istom uređaju. Kada se odredi veličina radne grupe, OpenCL deli veličinu globalnog posla na jedinice na uređaju. Radne jedinice u grupi se izvršavaju sinhronizovano i mogu da dele memoriju. Ukupan broj jedinica u radnoj grupi je nazvan veličina lokalnog posla. Grupne veličine se ne mogu dodeliti proizvoljno, OpenCL funkcija clGetKernelWorkGroupInfo() određuje veličinu grupe na uređaju.

Kada se odrede radne jedinice i veličina lokalnog posla, jezgro može biti poslato u komandni red za izvršavanje. Funkcija clEnqueueNDRangeKernel() omogućava postavljanje jezgra na uređaj zbog izvršavanja:

greska = clEnqueueNDRangeKernel( komandni_red, // ispravan komandi red jezgro,        //ispravan objekat jezgra 1,             // dimenzije problema posla NULL,          //rezervisano za buduće promene – mora biti NULL &global,       //radne jedinice za svaku dimenziju NULL,          // veličina grupnog posla za svaku dimenziju 0,             // broj događaja u listi događaja NULL,          // lista događaja koja treba da se kompletira pre izvršavanja NULL           // objekat događaja koji treba da se vrati na završetku ); 

Dimenzije problema mogu biti 1,2 i 3. Peti argument funkcije određuje veličinu radne jedinice za svaku dimenziju. Ako je potrebno obraditi sliku 512×512 pixela, mora biti određen niz koji pokazuje na broj radnih jedinica za svaku dimenziju:

size_t global[2]={512,512}; 

Šesti argument funkcije je broj radnih jedinica koje čine radnu grupu. Na primer, za 64 radnih jedinica koje su grupisane u 8×8 grupu, veličina za svaku dimenziju radne grupe se može odrediti na sledeći način:

size_t local[2]={8,8}; 

Ukupan broj radnih jedinica u svakoj grupi mora biti manji od maksimalnog broja radnih jedinica po grupi određenog funkcijom clGetKernelWorkGroupInfo(). Ako podaci ne moraju da se dele u radne grupe, postavljena je vrednost NULL i globalne radne jedinice se dele u odgovarajuće instance radnih grupa pri izvršavanju.

Sedmi i osmi argumenti kontrolišu niz izvršavanja komande jezgra. Ako karakteristike komandnog reda dozvoljavaju izvršavanje bez redosleda tokom clCreateCommandQueue() funkcije , lista događaja mora biti određena, zajedno s brojem događaja u listi koji treba da se kompletiraju pre nego što se određena komanda izvrši.Ako je lista prazna, broj elemnata u listi je nula. Poslednji argument dozvoljava instanci jezgra da kreira događaj na završetku. Ovo omogućava drugim komandama jezgra da čekaju na izvršavanje ove komande i da traže izvršavanje jezgra.

Ako je funkcija uspešno izvršena vraća se CL_SUCCESS.

Objekti memorije

OpenCL okvir obezbeđuje način da se paket podataka prevede u objekte memorije. Korišćenjem memorijskih objekata, OpenCL dozvoljava lako pakovanje svih podataka i lak prenos u memoriju uređaja tako da izvršavanje jezgra na uređaju ima lokalni pristup podacima.

Korišćenjem memorijskih objekata, minimizuju se memorijski prenosi od hosta i uređaja dok jezgro obrađuje podatke. OpenCL memorijski objekti se dele u dve grupe: objekte bafera i objekte slika. U objekte bafera se smeštaju jednodimenzioni podaci kao čto su int, float,vector… Objekti slika služe za smeštanje 2-dimenzionalnih i 3-dimenzionalnih nizova.

cl_mem ulaz; ulaz = clCreateBuffer( kontekst, // ispravan kontekst CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, // fleg bit polja koje određuje kako se koristi memorija sizeof(float) * VELICINA_PODATAKA, // veličina alociranog bafera u bajtovima podaci, // pokazivač na podatke bafera, koji treba da se iskopiraju s hosta &greska // greška ); 

Fleg bit polja se koristi za određivanje alokcije. Fleg definiše kako se memorija koristi. Neke od vrednosti flaga mogu biti:

• CL_MEM_READ_WRITE – jezgro može da piše i čita memoriju
• CL_MEM_WRITE_ONLY – jezgro može da piše
• CL_MEM_READ_ONLY – jezgro može da čita memoriju
• CL_MEM_COPY_HOST_PTR – OpenCL alocira memoriju i kopira podatke na koje pokazuje host_ptr u memorijski objekat

Ako je objekat bafera pravilno kreiran funkcija vraća objekat bafera s greškom čija je vrednost CL_SUCESS.

Kada se kreira memorijski objekat, komande se mogu smestiti u red da pišu podatke u objekat bafera iz memorije hosta ili da pročitaju podatke iz objakta bafera u memoriju hosta. OpenCL omogućava funkcije clEnqueueWriteBuffer() i clEnqueueReadBuffer() u ove svrhe.

greska = clEnqueueReadBuffer( komandi_red,            // ispravan komandi red izlazniBafer,                   // memorijski bafer iz kojeg se čita CL_TRUE,                  //indikator blokiranog čitanja 0,                        // odeljak od kojeg počinje čitanje u baferu sizeof(float) *DATA_SIZE, // veličina podataka koja treba da bude pročitana u bajtovima rezultat,                  // pokazivač u baferu na memoriji hosta gde treba da budu smešteni podaci 0,                        // broj događaja u listi događaja NULL,                     // lista događaja koja treba da se kompletira pre završetka NULL                      // objekat događaja koji se vraća pri završetku ); err = clEnqueueWriteBuffer( komandni_red, // ispravan red za komande ulazniBafer,         // objekat memorijskog bafera u koji se piše CL_TRUE,       //indikator blokiranog čitanja 0,             // odeljak od kojeg počinje pisanje u baferu sizeof(float) *DATA_SIZE, // veličina podataka koja treba da bude upisana u bajtovima host_ptr,      // pokazivač na bafer u memoriji hosta odakle treba da se pročitaju podaci 0,             // broj događaja u listi događaja NULL,          // lista događaja koja treba da se kompletira pre izvršavanja NULL           // objekat događaja koji se vraća pri izvršavanju ); 

Funkcija clEnqueueWriteBuffer() ređa komande koje pišu podatke iz memorije hosta u objekat bafera. Ova funkcija se obično koristi da obezbedi podatke za jezgro tokom procesiranja. Kao i kod funkcije clEnqueueReadBuffer() blokirajući fleg može biti CL_TRUE ili CL_FALSE da odredi da li je komanda za pisanje blokirajuća ili neblokirajuća.

Ako se uspešno završi funkcija vraća CL_SUCCESS.

Primeri

Primer 1: Program koji ispisuje „HelloWorld“

Jezgro pozdrav_svete_kernel.cl: __kernel void pozdrav_svete(__global char* ulazi, __global char* izlazi) { int broj = get_global_id(0); izlazi[broj] = ulazi[broj] + 1; } 

Program HelloWorld.cpp:

#include  #include  #include  #include  #include  #include  #include  #define USPESNO 0 #define NEUSPESNO 1 using namespace std; /* konvertovanje kernel fajla u string */ int konvertovanjeUString(const char *fajl, std::string& s) { size_t duzina; char*  str; std::fstream f(fajl, (std::fstream::in | std::fstream::binary)); if(f.is_open()) { size_t duzinaFajla; f.seekg(0, std::fstream::end); duzina = duzinaFajla = (size_t)f.tellg(); f.seekg(0, std::fstream::beg); str = new char[duzina+1]; if(!str) { f.close(); return 0; } f.read(str, duzinaFajla); f.close(); str[duzina] = '\0'; s = str; delete[] str; return 0; } cout<<"Greska: Fajl nije otvoren.\n:"< 0) { cl_platform_id* platforme = (cl_platform_id* )malloc(broj_platformi* sizeof(cl_platform_id)); status = clGetPlatformIDs(broj_platformi, platforme, NULL); platform = platforme[0]; free(platforme); } /*Zahtev za plaformom i odabir prvog GPU uredjaja ako takvog ima. Inace koristi CPU kao uredjaj.*/ cl_uint             broj_uredjaja = 0; cl_device_id        *uredjaji; status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, NULL, &broj_uredjaja); if (broj_uredjaja == 0) //nema dostupnih GPU uredjaja { cout < < "Nema dostupnih GPU uredjaja." << endl; cout << "Izaberi CPU podrazumenvani uredjaj." << endl; status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_CPU, 0, NULL, &broj_uredjaja); uredjaji = (cl_device_id*)malloc(broj_uredjaja * sizeof(cl_device_id)); status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_CPU, broj_uredjaja, uredjaji, NULL); } else { uredjaji = (cl_device_id*)malloc(broj_uredjaja * sizeof(cl_device_id)); status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, broj_uredjaja, uredjaji, NULL); } /*Kreiranje konteksta.*/ cl_context kontekst = clCreateContext(NULL,1, uredjaji,NULL,NULL,NULL); /*Kreiranje komandnog reda koji se povezuje s kontekstom*/ cl_command_queue komandi_red = clCreateCommandQueue(kontekst, uredjaji[0], 0, NULL); /*Kreiranje objekta programa */ const char *fajl = "pozdrav_svete_kernel.cl"; string kodString; status = konvertovanjeUString(fajl, kodString); const char *kod = kodString.c_str(); size_t duzina_koda[] = {strlen(kod)}; cl_program program = clCreateProgramWithSource(kontekst, 1, &kod, duzina_koda, NULL); /*Kreiranje programa. */ status=clBuildProgram(program, 1,uredjaji,NULL,NULL,NULL); /*Inicijalizacija ulaza,izlaza za host i kreiranje memorijskog objekta za jezgro*/ const char* ulaz = "GdkknVnqkc"; size_t str_duzina = strlen(ulaz); char *izlaz = (char*) malloc(str_duzina + 1); cl_mem ulazniBafer = clCreateBuffer(kontekst, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,          (str_duzina + 1) * sizeof(char),(void *) ulaz, NULL); cl_mem izlazniBafer = clCreateBuffer(kontekst, CL_MEM_WRITE_ONLY , (str_duzina + 1) * sizeof(char), NULL, NULL); /* Kreiranje objekat jezgra */ cl_kernel jezgro = clCreateKernel(program,"pozdrav_svete", NULL); /* Postavljanje argumenata jezgra*/ status = clSetKernelArg(jezgro, 0, sizeof(cl_mem), (void *)&ulazniBafer); status = clSetKernelArg(jezgro, 1, sizeof(cl_mem), (void *)&izlazniBafer); /* Pokretanje jezgra*/ size_t velicina_globalnog_posla[1] = {str_duzina}; status = clEnqueueNDRangeKernel(komandi_red, jezgro, 1, NULL, velicina_globalnog_posla, NULL, 0, NULL, NULL); /* Vracanje rezultata u memoriju hosta*/ status = clEnqueueReadBuffer(komandi_red, izlazniBafer, CL_TRUE, 0, str_duzina * sizeof(char), izlaz, 0, NULL, NULL); izlaz[str_duzina] = '\0'; cout << izlaz << endl; /*Oslobadjanje resursa:*/ status = clReleaseKernel(jezgro);               //Osobadjajnje jezgra status = clReleaseProgram(program);     //Osobadjajnje objekta programa status = clReleaseMemObject(ulazniBafer);       //Osobadjajnje memorijskog objekta status = clReleaseMemObject(izlazniBafer); status = clReleaseCommandQueue(komandi_red);    //Osobadjajnje komandnog reda status = clReleaseContext(kontekst);            //Osobadjajnje konteksta if (izlaz != NULL) { free(izlaz); izlaz = NULL; } if (uredjaji != NULL) { free(uredjaji); uredjaji = NULL; } std::string unos; std::getline (std::cin,unos); return USPESNO; } 

Rezultat izvrsavanja programa primera 1 je ispis „HelloWorld“, kako je prikazano na slici 3.

Primer 2: U ovom primeru preračunavanje približne vrednosti broja π se vrši tako što svaka radna jedinica generiše po 1000 tačaka i proverava da li tačke pripadaju jediničnom krugu. Jezgro koristi petostruki rekurzivni generator slučajnih brojeva za generisanje koordinata tačaka. Programu hosta se vraća niz čiji je svaki član broj tačaka koje pripadaju krugu iz jedne radne jedinice. Program host-a sabira sve članove niza, zatim preračunava broj π po formuli: π = 4*((broj_pogodaka * 1.0)/BROJ_TACAKA.

Jezgro izracunavanje_broja_Pi.cl: __kernel void izracunavanje_broja_Pi(__global int *izlaz) { long x[10]; long brUzoraka_u_jezgru=1000; long MAX = 4294967296; long randdmax = 2147483646; long m = 2147483647; int id = get_global_id(0); izlaz[id]=0; long s0 = id; long s1 = (69069*s0)%MAX; long s2 = (69069*s1)%MAX; long s3 = (69069*s2)%MAX; long s4 = (69069*s3)%MAX; long s5 = (69069*s4)%MAX; x[0]=s0%m; x[1]=s1%m; x[2]=s2%m; x[3]=s3%m; x[4]=s4%m; long a_1=107374182; long a_2=104480; long a_3=104480; long a_4=104480; long a_5=104480; double xx1; double yy1; double r; long z3; long z4; int i; i=5; int j; for(i=5;i<2*brUzoraka_u_jezgru+5;i+=2){ x[5]=(a_1*x[4]+a_5*x[0])%m; z3=x[5]; x[6]=(a_1*x[5]+a_5*x[1])%m; z4=x[6]; xx1=((1.0)*z3)/randdmax; yy1=((1.0)*z4)/randdmax; xx1=xx1*xx1; yy1=yy1*yy1; r=xx1+yy1; if(r <= 1) izlaz[id]=izlaz[id]+1; x[0]=x[1]; x[1]=x[2]; x[2]=x[3]; x[3]=x[4]; x[4]=x[5]; } }

Program HelloWorld.cpp:

#include 
#include 
#include 
#include 
#include 
#include 
#include 
#include 
#include 
#include 
#include 
#include  #include 
#define USPESNO 0
#define NEUSPESNO 1
#define BROJ_UZORAKA 120000
#define BROJ_PO_JEZGRU 1000
#define CLOCKS_PER_SEC 1000 using namespace std; /* konvertovanje kernel fajla u string */
int konvertovanjeUString(const char *fajl, std::string& s)
{
 size_t duzina;
 char* str;
 std::fstream f(fajl, (std::fstream::in | std::fstream::binary)); if(f.is_open())
 {
 size_t duzinaFajla;
 f.seekg(0, std::fstream::end);
 duzina = duzinaFajla = (size_t)f.tellg();
 f.seekg(0, std::fstream::beg);
 str = new char[duzina+1];
 if(!str)
 {
 f.close();
 return 0;
 } f.read(str, duzinaFajla);
 f.close();
 str[duzina] = '\0';
 s = str;
 delete[] str;
 return 0;
 }
 cout<<"Greska: Fajl nije otvoren.\n:"< return NEUSPESNO;
} int main(void) {
 const clock_t pocetni_trenutak = clock();
 cl_command_queue red_za_komande;
 cl_context kontekst;
 cl_device_id uredjaj;
 const cl_int broj_jezgara=BROJ_UZORAKA/BROJ_PO_JEZGRU;
 cl_int prvi_niz[broj_jezgara];
 long i;
 const size_t globalna_velicina_posla = sizeof(prvi_niz) / sizeof(cl_int);
 cl_kernel kernel;
 cl_mem bafer1;
 cl_platform_id platforma;
 cl_program program; clGetPlatformIDs(1, &platforma, NULL);
 clGetDeviceIDs(platforma, CL_DEVICE_TYPE_ALL, 1, &uredjaj, NULL);
 kontekst = clCreateContext(NULL, 1, &uredjaj, NULL, NULL, NULL);
 red_za_komande = clCreateCommandQueue(kontekst, uredjaj, 0, NULL);
 bafer1 = clCreateBuffer(kontekst, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(prvi_niz), &prvi_niz, NULL); /*Kreiranje objekta programa */
 const char *fajl = "izracunavanje_broja_Pi.cl";
 string kodString;
 cl_int status = konvertovanjeUString(fajl, kodString);
 const char *kod = kodString.c_str();
 size_t duzina_koda[] = {strlen(kod)};
 program = clCreateProgramWithSource(kontekst, 1, &kod, duzina_koda, NULL); clBuildProgram(program, 1, &uredjaj, "", NULL, NULL);
 kernel = clCreateKernel(program, "izracunavanje_broja_Pi", NULL);
 clSetKernelArg(kernel, 0, sizeof(cl_mem), &bafer1);
 //clSetKernelArg(kernel, 1, sizeof(cl_mem), &buffer2);
 clEnqueueNDRangeKernel(red_za_komande, kernel, 1, NULL, &globalna_velicina_posla, NULL, 0, NULL, NULL);
 clFlush(red_za_komande);
 clFinish(red_za_komande);
 clEnqueueReadBuffer(red_za_komande, bafer1, CL_TRUE, 0, sizeof(prvi_niz), &prvi_niz, 0, NULL, NULL);
 cl_double zbir = 0;
 for(i=0;i zbir += prvi_niz[i];
 }
 std::cout << "Priblizna vrednost broja Pi je " << 4*((zbir*1.0)/BROJ_UZORAKA) <<" ."<< endl;
 std::cout << "Broj uzoraka je " << BROJ_UZORAKA <<" ."<< endl;
 std::cout << "Paralelna obrada se izvrsavala " ;
 std::cout << float( clock () - pocetni_trenutak ) / CLOCKS_PER_SEC;
 std::cout << " sekundi." << endl;
 std::string unos;
 std::getline (std::cin,unos);

 return 1;
}

Mogući rezultat je 3.14117, za 5.023 sekunde, za ukupno 120 000 uzoraka, za 1000 uzoraka po jezgru. Jedan od mogućih rezultata je prikazan na slici 4. Rezultati nisu uvek identični, već približno jednaki.

Primer 3: U ovom primeru vrši se preračunavanje približne vrednosti broja π, korišćenjem funkcije za generisanje slučajnih brojeva iz biblioteke stdlib.h u programu host-a. U programu host-a se generišu dva niza slučajnih brojeva, jedan namenjen za generisanje vrednosti x koordinata, drugi za generisanje vrednosti y koordinata tačaka. Nizovi se prosleđuju jezgru, tako da svaka radna jedinica dobije po jednu tačku. Svaka radna jedinica preračunava da li se tačka koju je dobila nalazi u krugu. Svaka radna jedinica u odgovarajući član prvog niza upisuje 1 ako je tačka koju je dobila pripada krugu, ili 0 ako ne pripada. Program host-a prebrojava tačke koje pripadaju krugu i preračunva broj π po formuli:

π = 4*((broj_pogodaka * 1.0)/BROJ_TACAKA.

__kernel void izracunavanje_broja_Pi(__global int *izlaz1,__global int *izlaz2) { long max = 32767; int id = get_global_id(0); double x1=((1.0)*izlaz1[id])/max; double y1=((1.0)*izlaz2[id])/max; x1=x1*x1; y1=y1*y1; double r=x1+y1; if(r <= 1) izlaz1[id]=1; else izlaz1[id]=0; }

#include  #include  #include  #include  #include  #include  #include  #include  #include  #include  #include  #include  #include  #define USPESNO 0 #define NEUSPESNO 1 #define BROJ_UZORAKA 120000 #define CLOCKS_PER_SEC 1000 using namespace std; /* konvertovanje kernel fajla u string */ int konvertovanjeUString(const char *fajl, std::string& s) { size_t duzina; char* str; std::fstream f(fajl, (std::fstream::in | std::fstream::binary)); if(f.is_open()) { size_t duzinaFajla; f.seekg(0, std::fstream::end); duzina = duzinaFajla = (size_t)f.tellg(); f.seekg(0, std::fstream::beg); str = new char[duzina+1]; if(!str) { f.close(); return 0; } f.read(str, duzinaFajla); f.close(); str[duzina] = '\0'; s = str; delete[] str; return 0; } cout<<"Greska: Fajl nije otvoren.\n:"<
Na slici 5 je prikazan rezultat programa primera 3.
 
Korisni linkovi
http://amd-dev.wpengine.netdna-cdn.com/wordpress/media/2013/01/Introduction_to_OpenCL_Programming-Training_Guide-201005.pdf
https://en.wikipedia.org/wiki/OpenCL
Autor: Natasa Antic
Podeli ovo:
Click to share on Twitter (Opens in new window)
Click to share on Facebook (Opens in new window)