OpenCL - OpenGL Interopt

Ostatnio poznawałem podstawy OpenCL i kłopot sprawiło mi połączenie z OpenGL tak aby w zapisać wynik obliczeń w sposób łatwy i wygodny do wyświetlenia. Przyznam się, że miałem problem ze znalezieniem informacji jak to powinno wyglądać, jednak ostatecznie korzystając z doświadczenia z CUDA udało mi się to przy pomocy Pixel Buffer Object.

Żeby w ogóle móc połączyć OpenCL z OpenGL należy najpierw dodać nagłówek <CL/cl_gl.h> oraz dodać dodatkowe parametry do tworzenia kontekstu.

//Tylko dla Windowsa (dla Linuksa i Maca wygląda to inaczej)
cl_context_properties properties[] = {
             CL_CONTEXT_PLATFORM, (cl_context_properties)platform, 
             CL_GL_CONTEXT_KHR, (cl_context_properties)wglGetCurrentContext(),
             CL_WGL_HDC_KHR, (cl_context_properties)wglGetCurrentDC(), 0 };

Potem należy stworzyć PB:

unsigned int num_texels = width * height;
unsigned int num_values = num_texels * 4;
unsigned int size_tex_data = sizeof(GLubyte) * num_values;
void *data = malloc(size_tex_data);

// create buffer object
glGenBuffers(1, pbo);
glBindBuffer(GL_ARRAY_BUFFER, *pbo);
glBufferData(GL_ARRAY_BUFFER, size_tex_data, data, GL_DYNAMIC_DRAW);
free(data);

glBindBuffer(GL_ARRAY_BUFFER, 0);

Oraz teksturę, która posłuży do wyświetlenia:

GLuint textureId;
glGenTextures(1, &textureId);
glBindTexture(GL_TEXTURE_2D, textureId);
glTexParameterf(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_LINEAR);
glTexParameterf(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR_MIPMAP_LINEAR);
glTexParameterf(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_EDGE);
glTexParameterf(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_EDGE);
glTexParameteri(GL_TEXTURE_2D, GL_GENERATE_MIPMAP, GL_FALSE); // automatic mipmap generation included in OpenGL v1.4
glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA8, 640, 480, 0, GL_RGBA, GL_UNSIGNED_BYTE, 0);
glBindTexture(GL_TEXTURE_2D, 0);

W tej chwili mamy  teksturę i PBO teraz wystarczy już tylko przekazać do OpenCL informacje o buforze:

cl_mem clPBO = clCreateFromGLBuffer(context, CL_MEM_WRITE_ONLY, pbo, &error);

Teraz przykład wykonania kernela operującego na buforze (przed jego wywołaniem trzeba zarezerwować a potem zwolnić bufor).

clEnqueueAcquireGLObjects(cl_queue, 1, &clPBO, 0, NULL, NULL);
clSetKernelArg(kernel, 0, sizeof(cl_mem), &clPBO);

clEnqueueNDRangeKernel(cl_queue, kernel, 2, NULL, workSize, NULL, NULL, NULL, NULL);

clEnqueueReleaseGLObjects(cl_queue, 1, &clPBO, 0, NULL, NULL);

Jeżeli chcemy wyświetlić zawartość bufora najpierw trzeba uaktualnić teksturę:

glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pbo);
glBindTexture(GL_TEXTURE_2D, textureId);
glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, width, height, GL_RGBA, GL_UNSIGNED_BYTE, 0);

//Tu następuje rysowanie tekstury

glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);

Na koniec dobrze po sobie posprzątać i tyle jeżeli chodzi o zapisywanie wyniku obliczeń OpenCL do tekstury. Oczywiście to najprawdopodobniej nie jest jedyny sposób jednak póki co innego działającego który nie wymaga kopiowania danych z karty do RAMu i z powrotem już jako tekstura nie znalazłem.