Графикалық процессорларды программалаудың ерекшелiктерi

ГРАФИКАЛЫҚ ПРОЦЕССОРЛАРДЫ ПРОГРАММАЛАУДЫҢ ЕРЕКШЕЛIКТЕРI

Балганов Д.С.

OpenCL - Open Computing Language- Apple-дің гетерогендi жүйелердегi ағынды программалауы үшiн бастапқы жасалған тiлi, CPU және GPU –дан тұрады, қазiр Intel, AMD және NVIDIA стандартты қолдайды. Сондықтан, OpenCL-ға программалау туралы әңгiмелегеніміз жөн және матрицалардың көбейту мысалы арқылы осы технологияны төменде қарастырамыз ([1-3]).

Осыдан Brook+ пен CUDA-ның бір-біріне ұқсастығын көруімізге болады — екі жабдық та GPGPU пайдалануды қамтамасыз ететін үш мүмкіндікті береді: 1) олар ядроның есебін шығаратын кейбір тілді анықтайды. 2) олар жедел жад пен негізгі жүйелік жадтың арасындағы мәліметтерді алмастыру механизмін көрсетеді. 3) олар әдеттегі C/C+—+ немесе FORTRAN (CUDA-мен болған жағдайда) бағдарламалау тілін ядромен орындалатын тізбекті іске қосу үшін синтаксистік құрылымдармен кеңейтеді.

OpenCL басқаша орналастырылған. Бұл жабдықтың бастапқы бастамасы ретінде Khronos Group шеңберіндегі Apple болған, ол OpenGLдің және графикті қарастырудың стандартының дамуымен жұмыс жасады. Бұл өз ізін қалдырды: OpenCL OpenGL сияқты қандайда бір бағдарламалау тілімен пайдалануға болатын тек кітапхана ретінде болады (мысалы, Python-ға OpenCL-дің байланысы бар, Ruby және т.б.). Әрине, Brook+ ға және C for CUDAға арналған бағдарламалар нәтижесінде кодқа трансталады, ол жұмысты орындау үшін кітапханалық қызметтерді қабылдау ретінде ғана пайдаланылады, бұл қызметтерді қабылдау CUDA үшін және ATI/AMD Stream үшін де белгілі, оларды, әрине, бағдарламада пайдалануға болады. Бірақ OpenCL қызметтерді таңдау арқылы ғана есептік ресурстарды пайдалануды ұсынады.

Сонымен қатар, әрине, OpenCL C ядросының сипаты үшін тілді де анықтайды. Бірақ тек ядроны бағдарламалау амалы ретінде ғана. Бұл OpenCL C үшін жеке хабарлау көрсетілмеген, орындалатын машиналық кодта OpenCL C-ға шығыс мәтіні кітапхананың арнайы қызметіне айналады.

Мүмкін, бұл жерде матрицалар қосындысына қарапайым мысал келтіре салудың қажеті жоқ, ядроның өзіндік ерекшелігі анық. OpenCL үшін бағдарламалар көлемді, уақыт пен орынды экономдау үшін матрицаны көбейту керек. Төменде сәйкесінше есептеуді шығаратын негізгі бағдарлама мен матрицаны көбейтудің ядросы көрсетілген.

OpenCL-ді бағдарламалаудың барлық тілімен пайдалануға болады. Төменде көрсетілген бағдарлама өзінің орындалуымен (теориялық) CUDA ға көрсетілген бағдарламаға жақын. Ядроның коды да бірдей көрінеді (қаласаң, оны да көруге болады). Келесідей жағдайлар ерекшеленеді:

Локалды, жедел, барлық желіге рұқсат етілген, бір локалды топқа жиналған жад бар. Сәйкесінше,

kernel void matrix_multiply(

unsigned AH, unsigned AW, __global float* AE,

unsigned BH, unsigned BW, __global float* BE,

unsigned RH, unsigned RW, __global float* RE)

{

float sum = 0; int i = get_global_id(1);

for( k = 0; k < AW; k += 1) {

sum += AE[i * AW + k] * BE[k * BW + j];

}

RE [i * RW + j] = sum;

OpenCL. Матрицаларды көбейтудiң ядросы

түрлендiргiштер global және local. Матрицаны көбейтудің басты бағдарламасы бірінші қолданылған, көрсетiлу үшiн бiрiншi қолданған онда, матрица оқып отыратын және есептеудiң нәтижелерiн үдеткiштiң жаhандық жадына жазу керек.

Үшіншіден, есептеуді іске қосу үдерісі. Осыны іске қосу үшін орындалатын логикалық қадамдар Brook+ мен CUDA-ға да осындай қадамдар жасау керек. Бірақ техникалық, OpenCL кітапхана ретінде ғана қызмет атқарады, бағдарламалаудың негізгі тілін қандайда бір жаңа базалық түрлермен, синтаксистік құрылымдармен кеңейтпейді, желіні іске қосу үшін көп операция мен деталь жасау керек және олар Brook+ или C for CUDA-та сәйкес компилятормен автоматты түрде анықтайды және қалыптасады.

Матрицаны көбейтудің басты бағдарламасындағы кодты қарай отырыпOpenCL қызметін пайдалану толығымен қате коды әкелуі мүмкін, және бұл қателер жақсы бағдарламаларда жұмыс істете білу керек. Мысалда қайтадан олар келтірілмеген, себебі онсыз да үлкен (мәтін үшін) код ала алмайсыз, тіпті нашарлап кетеді.

Жеделдетуге бағыттаушының және негізгі жүйенің өлшемі әртүрлі болып келеді(мысалы, 32 және 64 бит), сондықтан жедел жад аймағында бағыттаушының мәні үшін негізгі жүйеде cl_mem түрі пайдаланылады. Ядроға бұндай мәнді жіберу кезінде OpenCL аргумент-бағыттаушы ретінде керекті қалыпқа ие болуы керек. Сондықтан 7-суретте қарастырылған мысалда ядродағы маррицалар жайлы ақпарат өзі бағыттаушы болып келетін жолда құрылым түрінде жіберіледі, ал OpenCL үшін мәліметтер түрі қарапайымдығы арқылы, мәліметтер одағына корректі бағыттаушыға cl_mem түрінің мәні сәйкес аргументтер үшін (олардың номерлері 2, 5, 8-ден басталады) іске қосу жүйесіне айналдыруды мақсат етеді.

Жеделдетумен жұмыс командалар кезегі мәтінінде бір құрылғының пайда болуынан және жұмысқа мәтіннің пайда болуынан пайда басталады (22-27 жолдар). Мұнда мәтін GPU үшін құрылады, термин ретінде OpenCL есептеумен қатар OpenGL или DirectX интерфейстері үшін графикамен жұмысты қамтамасыз ете алатын құрылғы дегенді білдіреді. Мәтінді басқа құрылғыларға да құрауға болады, мысалы, (CL_DEVIC_TYPE_CPU) орталық процессорлары үшін немесе Cell (CL_DEVICE_TYPE_ACCELERATOR) процессорлар негізінде жеделдетулер үшін немесе (CL_DEVICE_TYPE_ALL) жүйесінде барлық OpenCL құрылғылары үшін құралады. Аталғандай, OpenCL-ды бірегей құрылғы қылдырып жасағысы келеді.

Мәтін қандай да бинарлы кодқа ядроның шығыс мәтінін хабарлау үшін және қандай құрылғыларда желі іске қосылады екенін және т.б. анықтайды.

Ядро жасайтын желіні іске қосу топтары немесе жедел жадтан мәліметтер жадына баспаларды оқу операциясы міндетті түрде тәртіппен тұрған емес арнайы командалар кезегімен орындалады. Бұл командаларды аратұра жүйелеу үшін (events) жағдайымен, OpenCL анықтайтын операциялармен жүйелеуге болады. Сәйкесінше мыналар: командалар кезегі (command queue) есептеуді ол арқылы жіберуге болу үшін құрылады. Ол екінші құрылғыға жалғанады (қол жетімділер арасындағылардың біріншісіне).

Ары қарай (29-31 жолдар) бағдарламаның компиляциясы және одан ядроның пайда болуымен орындалады. Ядроның бастапқы мәтін бағдарламасы, байттар тізбегі жадқа жүктелген және оған program_source нұсқайды. Хабарлау бұрын құрылғаан мәтінде жасалады, яғни GPU класының құрылғылары үшін құрылған. Бағдарламаны бір ғана бастапқы мәтіннен емес, бірнеше (берілген мысалда clCreateProgramWithSource шақырудың екінші аргументінде бір мәтінді пайдалану көрсетілген) мәтіннен қалыптастыруға болады.

0. 
   typedef struct
   
1. 
   {
   
2. 
   unsigned W; unsigned H; cl_mem E;
   
3. 
   } matrix_t;
   
4. 
   
   
5. 
   int main(int argc, char** argv)
   
6. 
   {
   
7. 
   const unsigned L = 1234, M = 567, N = 8910;
   
8. 
   const unsigned BLOCK_SIZE = 16;
   
9. 
   cl_context cntx ;
   
10. 
    cl_command_queue cq;
    
11. 
    cl_device_id* ds; // devices
    
12. 
    cl_program p;
    
13. 
    cl_kernel k;
    
14. 
    char* program_source = "..."; // kernel source code
    
15. 
    size_t global_work_size [2] = {L , N};
    
16. 
    size_t local_work_size [2] = {BLOCK_SIZE, BLOCK_SIZE};
    
17. 
    size_t cb ;
    
18. 
    float mem_A[L] [M] , mem_B[M] [N] , mem_R[L][N];
    
19. 
    matrix_t A, B, R;
    
20. 
    int i = -1;
    
21. 
    // ...
    
22. 
    cntx = clCreateContextFromType(NULL , CL_DEVICE_TYPE_GPU , NULL, NULL, NULL);
    
23. 
    clGetContextInfo(cntx, CL_CONTEXT_DEVICES , 0, NULL, &cb);
    
24. 
    ds = malloc(cb);
    
25. 
    clGetContextInfo(cntx, CL_CONTEXT_DEVICES, cb, ds, NULL);
    
26. 
    cq = clCreateCommandQueue(cntx, ds[0], 0, NULL);
    
27. 
    free(ds);
    
28. 
    
    
29. 
    p = clCreateProgramWithSource(cntx, 1, &program_source, NULL, NULL);
    
30. 
    clBuildProgram(p, 0, NULL , NULL , NULL , NULL);
    
31. 
    k = clCreateKernel(p, "matrix_multiply", NULL);
    
32. 
    
    
33. 
    A.H = L; A.W = M;
    
34. 
    A.E = clCreateBuffer(cntx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
    
35. 
    A.H * A.W * sizeof(cl_float), mem_A, NULL);
    
36. 
    B.H = M; B.W = N;
    
37. 
    B.E = clCreateBuffer(cntx, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
    
38. 
    B.H * B.W * sizeof(cl_float), mem_B, NULL);
    
39. 
    R. H = L; R. W = N;
    
40. 
    R.E = clCreateBuffer(cntx, CL_MEM_READ_WRITE ,
    
41. 
    R.H * R.W * sizeof(cl_float), NULL, NULL);
    
42. 
    
    
43. 
    clSetKernelArg(k, i += 1, sizeof(unsigned), &A.H);
    
44. 
    clSetKernelArg(k, i += 1, sizeof(unsigned), &A.W);
    
45. 
    clSetKernelArg(k, i += 1, sizeof(cl_mem), &A.E);
    
46. 
    clSetKernelArg(k, i += 1, sizeof(unsigned), &B.H);
    
47. 
    clSetKernelArg(k, i += 1, sizeof(unsigned), &B.W);
    
48. 
    clSetKernelArg(k, i += 1, sizeof(cl_mem), &B.E);
    
49. 
    clSetKernelArg(k, i += 1, sizeof(unsigned), &R.H);
    
50. 
    clSetKernelArg(k, i += 1, sizeof(unsigned), &R.W);
    
51. 
    clSetKernelArg(k, i += 1, sizeof(cl_mem), &R.E);
    
52. 
    
    
53. 
    clEnqueueNDRangeKernel(cq, k, 2, NULL,
    
54. 
    global_work_size, local_work_size, 0, NULL, NULL);
    
55. 
    clEnqueueReadBuffer(cq, R.E, CL_TRUE, 0,
    
56. 
    R.W * R.H * sizeof(cl_float), mem_R, 0, NULL, NULL);
    
57. 
    
    
58. 
    clReleaseMemObject(A.E);
    
59. 
    clReleaseMemObject(B.E);
    
60. 
    clReleaseMemObject(R.E);
    
61. 
    clReleaseKernel( k) ;
    
62. 
    clReleaseProgram( p) ;
    
63. 
    clReleaseCommandQueue(cq);
    
64. 
    clReleaseContext(cntx);
    
65. 
    // ...
    
66. 
    return 0;
    
67. 
    }
    

Келесі әрекеттер (33-41 жолдар) жедел жадпен ерекшеленеді. А және В матрицалары үшін шығыс мәліметтері ретінде онымен бір уақытта толтырылады. OpenCL-да бұны бір шақырумен орындауға болады әрі бұл өте тиімді.

Содан кейін ядро үшін аргументтер орнататын код одағы келеді (43-51 жолдар). Бұнда cl_mem түрі мәнінде жедел жад аймағында осындай мәнді қайталауға болады. OpenCL орындау жүйесі ядроның кейбір аргументтері бағыттаушы екенін анықтаған кезде ол бұл бағыттаушыны дұрыс қалыптастыру үшін cl_mem түрінің мәнін жіберуге әрeкет етеді. Сәйкесінше, теориялық, OpenCL ядросына жад аймағы жайлы ақпаратты басқа тәсілмен жіберуге болмайды.

Сонымен қатар, ядроға рагумент керекті мәліметтер өлшемі (clSetKernelArg үшінші аргументі) көшірілу арқылы қарапайым түрде қажетті адрес бойынша белгіленіп жіберіледі (4-аргумент).

(53-56 жолдар) есептеу мен оның нәтижесін алуды жіберу. Бұл кезекке командалардың жазылуы арқылы орындалады. Бірінші команда NDRangeKernel (ND — N-Dimension), бұл ядро жіберу командасы болып келеді. Онда индекстің (2, global_work_size —L X N өлшемді екідеңгейлік аймақ) жаһанды кеңістігінің өлшемі көрсетіледі, әрқайсысына (local_work_size —CUDAға мысалдағындай одақ өлшемі : 16 X 16) локалды желілер топтарына өлшемді желілер жіберіледі.

Орнату кезекке қойғаннан кейін басқарма негізгі бағдарламаға қайтып келеді, есептеу негізгі лекке сәйкес басқару асинхронды болады.

Кезекке келген келесі команда — ReadBuffer. Оның мағынасы түсінікті, ал бар көңілді CL_TRUE ала алады. Яғни OpenCL оқудың осы операциясын күту керек, яғни қалай басқарманы негізгі басдарламаға қайтаруды ойлау керек.

Сонымен қатар, мысалда кезекке тұрған екі команда да жағдай жайлы ешқандай ақпаратсыз жазылады (соңғы үш параметр шақыруда нөлге тең: 0, NULL, NULL). OpenCL үшін бұл бұл командаларды олардың кезекке тұру реті бойынша орындау керек дегенді білдіреді. Яғни, ReadBuffer NDRangeKernel-дан кейін орындалады, одан кейін басқару негізгі бағдарламаға қайтып келеді, содан mem_R-ден мәліметтерді пайдалануға болады. Осыны дәлелдеу керек еді.

OpenCL үшін double түрінің мәні жайлы сұрақтар қалды. Жалпы, осындай мәнділермен жұмыс жасау стандартты қажет етпейді, бірақ олармен жұмыс жасауға болады. NVIDIA и AMD-дан OpenCL 1.0 ағымдағы жүзеге асырулар осындай түрде қолданыла алмайды.

1.By M., Девис Т., Нейдер Дж., Шрайнер Д. Open GL. Руководство по программированию. Библиотека программиста. 4-e издание. –СПб.: Питер, 2006.-624c.: ил.