OpenCL

OpenCL é um padrão aberto e livre de royalties para a programação paralela em ambientes computacionais heterogêneos, compostos de CPUs, GPUs e outros processadores. OpenCL permite a escrita de código multi-plataforma para execução nestes dispositivos, possibilitando a utilização de todo o poder computacional disponível no ambiente.

Para mais informações sobre o padrão, consulte o site do Khronos Group: http://www.khronos.org/opencl/

Minicursos

A V3D vem apresentando minicursos de OpenCL em alguns eventos no país. Abaixo, você pode realizar o download do material disponibilizado nestes minicursos.

Programação introdutória em OpenCL e aplicações em realidade virtual e aumentada
César L. B. Silveira e Luiz G. da Silveira Jr.
SVR 2010 (XII Symposium on Virtual and Augmented Reality)
[Texto (PDF)] [Slides (PDF)]

Programação Paralela em Ambientes Computacionais Heterogêneos com OpenCL
César L. B. Silveira, Luiz G. da Silveira Jr e Gerson Geraldo H. Cavalheiro
WSCAD 2010 (XI Simpósio em Sistemas Computacionais)
[Texto (PDF)] [Slides (PDF)] [Slides (PPTX)]

Programação em OpenCL: Uma introdução prática
César L. B. Silveira, Luiz G. da Silveira Jr e Gerson Geraldo H. Cavalheiro
SBGames 2010 (IX Simpósio Brasileiro de Jogos e Entretenimento Digital)
[Texto (PDF)] [Slides (PDF)] [Slides (PPTX)]

Programação paralela em OpenCL
César L. B. Silveira
GDS 2010 (Game Development School Unisinos 2010)
[Slides (PDF)] [Slides (PPTX)]

Exemplos de código

Subtração de elementos de dois arrays

O exemplo a seguir demonstra a configuração de um contexto OpenCL para a execução de um kernel cuja finalidade é realizar a subtração dos elementos de dois array de entrada, armazenando o resultado em um terceiro array. Os arrays são preenchidos com valores aleatórios. Ao final, os dados de entrada e saída são impressos no console.

#include <stdio.h>
#include <stdlib.h>
 
#ifdef __APPLE__
#include <OpenCL/opencl.h>
#else
#include <CL/opencl.h>
#endif
 
#define ARRAY_LENGTH 1000
 
int main(int argc, char** argv)
{
    /* Variáveis para armazenamento de referências a
       objetos OpenCL */
    cl_platform_id platformId;
    cl_device_id deviceId;
    cl_context context;
    cl_command_queue queue;
    cl_program program;
    cl_kernel kernel;
    cl_mem bufA;
    cl_mem bufB;
    cl_mem bufC;
 
    /* Variáveis diversas da aplicação */
    int* hostA;
    int* hostB;
    int* hostC;
    size_t globalSize[1] = { ARRAY_LENGTH };
    int i;
 
    /* Código-fonte do kernel */
    const char* source =
    "__kernel void ArrayDiff( \
        __global const int* a, \
        __global const int* b, \
        __global int* c) \
    { \
        int id = get_global_id(0); \
        c[id] = a[id] - b[id]; \
    }";
 
    /* Obtenção de identificadores de plataforma
       e dispositivo. Será solicitada uma GPU. */
    clGetPlatformIDs(1, &platformId, NULL);
    clGetDeviceIDs(platformId, CL_DEVICE_TYPE_GPU,
        1, &deviceId, NULL);
 
    /* Criação do contexto */
    context = clCreateContext(0, 1, &deviceId,
        NULL, NULL, NULL);
 
    /* Criação da fila de comandos para o
       dispositivo encontrado */
    queue = clCreateCommandQueue(context, deviceId,
        0, NULL);
 
    /* Criação do objeto de programa a partir do
       código-fonte armazenado na string source */
    program = clCreateProgramWithSource(context, 1, &source,
        NULL, NULL);
 
    /* Compilação do programa para todos os
        dispositivos do contexto */
    clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
 
    /* Obtenção de um kernel a partir do
        programa compilado */
    kernel = clCreateKernel(program, "ArrayDiff", NULL);
 
    /* Alocação e inicialização dos arrays no hospedeiro */
    hostA = (int*) malloc(ARRAY_LENGTH * sizeof(int));
    hostB = (int*) malloc(ARRAY_LENGTH * sizeof(int));
    hostC = (int*) malloc(ARRAY_LENGTH * sizeof(int));
 
    for (i = 0; i < ARRAY_LENGTH; ++i)
    {
        hostA[i] = rand() % 101 - 50;
        hostB[i] = rand() % 101 - 50;
    }
 
    /* Criação dos objetos de memória para comunicação com
       a memória global do dispositivo encontrado */
    bufA = clCreateBuffer(context, CL_MEM_READ_ONLY,
        ARRAY_LENGTH * sizeof(int), NULL, NULL);
 
    bufB = clCreateBuffer(context, CL_MEM_READ_ONLY,
        ARRAY_LENGTH * sizeof(int), NULL, NULL);
 
    bufC = clCreateBuffer(context, CL_MEM_READ_WRITE,
        ARRAY_LENGTH * sizeof(int), NULL, NULL);
 
    /* Transferência dos arrays de entrada para a memória
       do dispositivo */
    clEnqueueWriteBuffer(queue, bufA, CL_TRUE, 0,
        ARRAY_LENGTH * sizeof(int), hostA, 0,
        NULL, NULL);
 
    clEnqueueWriteBuffer(queue, bufB, CL_TRUE,  0,
        ARRAY_LENGTH * sizeof(int), hostB, 0,
        NULL, NULL);
 
    /* Configuração dos argumentos do kernel */
    clSetKernelArg(kernel, 0, sizeof(cl_mem), &bufA);
    clSetKernelArg(kernel, 1, sizeof(cl_mem), &bufB);
    clSetKernelArg(kernel, 2, sizeof(cl_mem), &bufC);
 
    /* Envio do kernel para execução no dispositivo */
    clEnqueueNDRangeKernel(queue, kernel, 1, NULL,
        globalSize, NULL, 0, NULL, NULL);
 
    /* Sincronização (bloqueia hospedeiro até término da
       execução do kernel */
    clFinish(queue);
 
    /* Transferência dos resultados da computação para a
       memória do hospedeiro */
    clEnqueueReadBuffer(queue, bufC, CL_TRUE, 0,
        ARRAY_LENGTH * sizeof(int), hostC, 0,
        NULL, NULL);
 
    /* Impressão dos resultados na saída padrão */
    for (i = 0; i < ARRAY_LENGTH; ++i)
    {
        printf("%d - %d = %d\n", hostA[i], hostB[i], hostC[i]);
    }
 
    /* Liberação de recursos e encerramento da aplicação */
    clReleaseMemObject(bufA);
    clReleaseMemObject(bufB);
    clReleaseMemObject(bufC);
    clReleaseKernel(kernel);
    clReleaseProgram(program);
    clReleaseCommandQueue(queue);
    clReleaseContext(context);
 
    free(hostA);
    free(hostB);
    free(hostC);
 
    return 0;
}

Interoperação com OpenGL

O exemplo de código a seguir demonstra a possibilidade de interoperação entre OpenCL e OpenGL. Quando executado, este código exibe uma malha tridimensional animada segundo uma função senóide. Os vértices da malha são armazenados em um Vertex Buffer Object (VBO) no contexto OpenGL. Este VBO é compartilhado com o contexto OpenCL, no qual é executado um kernel responsável pelo posicionamento dos vértices da malha de acordo com o tempo da animação.

#include <stdio.h>
#include <stdlib.h>
#include <GL/glew.h>
#include <GL/gl.h>
#include <GL/glut.h>
#include <CL/opencl.h>
 
#ifndef _WIN32
#include <GL/glx.h>
#endif
 
/* Dimensões da janela e da malha */
const unsigned int windowWidth = 512;
const unsigned int windowHeight = 512;
const unsigned int meshWidth = 256;
const unsigned int meshHeight = 256;
 
/* Objetos OpenGL */
GLuint vbo;
int window = 0;
 
/* Objetos OpenCL */
cl_platform_id platformId;
cl_context context;
cl_device_id deviceId ;
cl_command_queue queue;
cl_kernel kernel;
cl_mem vboCL;
cl_program program;
 
/* Código-fonte do kernel */
const char* kernelSource = " \
__kernel void sine_wave( \n  \
   __global float4* pos, \n  \
   unsigned int width, \n    \
   unsigned int height, \n   \
   float time) \n            \
{ \n                                            \
   unsigned int x = get_global_id(0); \n        \
   unsigned int y = get_global_id(1); \n        \
\n                                              \
   float u = x / (float) width; \n              \
   float v = y / (float) height; \n             \
   u = u*2.0f - 1.0f; \n                        \
   v = v*2.0f - 1.0f; \n                        \
\n                                              \
   float freq = 4.0f; \n                                             \
   float w = sin(u * freq + time) * cos(v * freq + time) * 0.5f; \n  \
\n                                                                   \
   pos[y * width + x] = (float4)(u, w, v, 1.0f); \n                  \
}";
 
/* Tempo de simulação (determina posição dos vértices na malha) */
float anim = 0.0;
 
void initGL(int argc, char** argv);
void initCL();
void displayGL(void);
void keyboard(unsigned char key, int x, int y);
void cleanup();
 
int main(int argc, char** argv)
{
   /* Inicialização GLUT */
   glutInit(&argc, argv);
   glutInitDisplayMode(GLUT_RGBA | GLUT_DOUBLE);
   glutInitWindowPosition(glutGet(GLUT_SCREEN_WIDTH)/2 - windowWidth/2,
                          glutGet(GLUT_SCREEN_HEIGHT)/2 - windowHeight/2);
   glutInitWindowSize(windowWidth, windowHeight);
   window = glutCreateWindow("OpenCL/GL Interop (VBO)");
   glutDisplayFunc(displayGL);
   glutKeyboardFunc(keyboard);
 
   /* Inicialização OpenGL e OpenCL*/
   initGL(argc, argv);
   initCL();
 
   glutMainLoop();
 
   cleanup();
 
   return 0;
}
 
void initGL(int argc, char** argv)
{
   glClearColor(0.0, 0.0, 0.0, 1.0);
   glDisable(GL_DEPTH_TEST);
   glViewport(0, 0, windowWidth, windowHeight);
   glMatrixMode(GL_PROJECTION);
   glLoadIdentity();
   gluPerspective(
      60.0, (GLfloat) windowWidth / (GLfloat) windowHeight, 0.1, 10.0);
   glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);
   glMatrixMode(GL_MODELVIEW);
   glLoadIdentity();
   glTranslatef(0.0, 0.25, -3.0);
   glRotatef(45.0, 1.0, 0.0, 0.0);
   glRotatef(45.0, 0.0, 1.0, 0.0);
 
   glewInit();
 
   /* Criação do VBO */
   glGenBuffers(1, &vbo);
   glBindBuffer(GL_ARRAY_BUFFER, vbo);
   glBufferData(
      GL_ARRAY_BUFFER, meshWidth * meshHeight * 4 * sizeof(float),
      0, GL_DYNAMIC_DRAW);
}
 
void initCL()
{
   /* Obtenção de identificadores de plataforma e dispositivo (GPU) */
   clGetPlatformIDs(1, &platformId, NULL);
   clGetDeviceIDs(platformId, CL_DEVICE_TYPE_GPU, 1, &deviceId, NULL);
 
   /* Criação do contexto com propriedades para o compartilhamento
    * com OpenGL */
#ifdef _WIN32
   cl_context_properties props[] =
      {
         CL_GL_CONTEXT_KHR, (cl_context_properties) wglGetCurrentContext(),
         CL_WGL_HDC_KHR, (cl_context_properties) wglGetCurrentDC(),
         CL_CONTEXT_PLATFORM, (cl_context_properties) platformId,
         0
	  };
#else
   cl_context_properties props[] =
      {
         CL_GL_CONTEXT_KHR, (cl_context_properties)glXGetCurrentContext(),
         CL_GLX_DISPLAY_KHR, (cl_context_properties)glXGetCurrentDisplay(),
         CL_CONTEXT_PLATFORM, (cl_context_properties)platformId,
         0
      };
#endif
   context = clCreateContext(props, 1, &deviceId, NULL, NULL, NULL);
 
   /* Criação da fila de comandos do dispositivo */
   queue = clCreateCommandQueue(context, deviceId, 0, NULL);
 
   /* Criação de compilação do programa e do kernel */
   program = clCreateProgramWithSource(
      context, 1, (const char **) &kernelSource, NULL, NULL);
   clBuildProgram(program, 0, NULL, "-cl-fast-relaxed-math", NULL, NULL);
   kernel = clCreateKernel(program, "sine_wave", NULL);
 
   /* Criação do objeto de memória para o VBO OpenGL */
   vboCL = clCreateFromGLBuffer(context, CL_MEM_WRITE_ONLY, vbo, NULL);
 
   /* Configuração dos argumentos fixos do kernel */
   clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &vboCL);
   clSetKernelArg(kernel, 1, sizeof(unsigned int), &meshWidth);
   clSetKernelArg(kernel, 2, sizeof(unsigned int), &meshHeight);
}
 
void displayGL()
{
   /* Dimensões do espaço índices */
   const size_t globalSize[2] = { meshWidth, meshHeight };
 
   /* Incremento do tempo de simulação */
   anim += 0.01f;
 
   /* Aguarda término das operações OpenGL */
   glFinish();
 
   /* Aquisição do VBO para uso com OpenCL */
   clEnqueueAcquireGLObjects(queue, 1, &vboCL, 0,0,0);
 
   /* Atualização do tempo de simulação no kernel */
   clSetKernelArg(kernel, 3, sizeof(float), &anim);
 
   /* Execução do kernel */
   clEnqueueNDRangeKernel(queue, kernel, 2, NULL, globalSize, NULL, 0,0,0 );
 
   /* Liberação do VBO de volta para OpenGL */
   clEnqueueReleaseGLObjects(queue, 1, &vboCL, 0,0,0);
 
   /* Sincronização */
   clFinish(queue);
 
   /* Renderização da malha */
   glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);
   glBindBuffer(GL_ARRAY_BUFFER, vbo);
   glVertexPointer(4, GL_FLOAT, 0, 0);
   glEnableClientState(GL_VERTEX_ARRAY);
   glColor3f(1.0, 0.0, 0.0);
   glDrawArrays(GL_POINTS, 0, meshWidth * meshHeight);
   glDisableClientState(GL_VERTEX_ARRAY);
   glutSwapBuffers();
   glutPostRedisplay();
}
 
void keyboard(unsigned char key, int x, int y)
{
   if (key == 'q' || key == 'Q')
	   cleanup();
}
 
void cleanup()
{
   clReleaseKernel(kernel);
   clReleaseProgram(program);
   clReleaseCommandQueue(queue);
 
   glBindBuffer(1, vbo);
   glDeleteBuffers(1, &vbo);
   clReleaseMemObject(vboCL);
 
   clReleaseContext(context);
   glutDestroyWindow(window);
 
   exit(0);
}

Toolkits OpenCL

Para desenvolver soluções OpenCL, é necessário instalar um toolkit de desenvolvimento OpenCL para os dispositivos que você deseja utilizar. Seguem links para alguns destes toolkits:

No caso de GPUs, também é necessário realizar o download e a instalação de drivers de desenvolvimento.

V3D Labs is proudly powered by WordPress and Siteslike