OpenCL

 Nota: Não confundir com OpenGL.

OpenCL (Open Computing Language) é uma arquitetura para escrever programas que funcionam em plataformas heterogêneas, consistindo em CPUs, GPUs e outros processadores. OpenCL inclui uma linguagem (baseada em C99) para escrever kernels (funções executadas em dispositivos OpenCL), além de APIs que são usadas para definir e depois controlar as plataformas heterogênea. OpenCL permite programação paralela usando, tanto o paralelismo de tarefas, como de dados.

Imagem ilustrativa

Ela foi adotada para controladores de placas gráficas pela AMD/ATI, que a tornou na sua única oferta de GPU como Stream SDK, e pela Nvidia, que oferece também OpenCL como a escolha para o seu Compute Unified Device Architecture (CUDA) nos seus controladores.

A arquitetura OpenCL partilha uma série de interfaces computacionais, tanto com CUDA, como com a concorrente DirectCompute da Microsoft.

A proposta OpenCL é similar às propostas OpenGL e OpenAL, que são padrões abertos da indústria para gráficos 3D e áudio, respectivamente. OpenCL estende o poder da GPU além do uso gráfico (GPGPU). OpenCL é gerido pelo consórcio tecnológico Khronos Group.

Histórico

OpenCL foi inicialmente desenvolvida pela Apple Inc., que tem direitos sobre a marca, e produziu uma proposta inicial em colaboração com equipes técnicas na AMD, IBM, Intel e Nvidia. Apple submeteu esta proposta inicial ao Khronos Group. Em 16 de junho de 2008, o Khronos Compute Working Group foi formado com representantes de companhias fabricantes de CPU, GPU, embedded-processor e software. Este grupo trabalhou por cinco meses para finalizar os detalhes técnicos da especificação OpenCL 1.0 em 18 de novembro de 2008.

Esta especificação técnica foi revisada por membros da Khronos e aprovada como release pública em 8 de dezembro de 2008.

OpenCL 1.0 foi introduzida no Mac OS X v10.6 ('Snow Leopard') em 28 de agosto de 2008. De acordo com a nota de imprensa:

Snow Leopard terá extensões para suporte à hardware moderno com Open Computing Language (OpenCL), que permite qualquer aplicação explorar os vastos gigaflops do poder de computação da GPU antes disponíveis somente para aplicações gráficas. OpenCL é baseado na linguagem de programação C e está sendo proposto como um padrão aberto.

AMD decidiu oferecer suporte a OpenCL (e DirectX 11) ao invés do agora obsoleto Close to Metal em seu Stream framework.

RapidMind anunciou a adoção de OpenCL sob sua plataforma de desenvolvimento, para permitir suporte a GPUs de múltiplos fabricantes com uma única interface.

Nvidia anunciou em 9 de dezembro de 2008 agregar suporte completo à especificação OpenCL 1.0 ao seu GPU Computing Toolkit.

Nvidia anunciou em 13 de maio de 2009 aos desenvolvedores o "OpenCL 1.0 Conformance Candidate Release", é a primeira solução de compilação de códigos desenvolvidos em OpenCL.

Exemplo

O modelo de programação OpenCL é bastante semelhante ao modelo de programação CUDA de Nvidia. Abaixo temos os códigos fontes de dois núcleos que fazem a adição de dois vetores:

Código C para um núcleo CUDA:

  __global__ void
  vectorAdd(const float * a, const float * b, float * c)
  {
      // Indice do vetor
      int nIndex = blockIdx.x * blockDim.x + threadIdx.x;
      c[nIndex] = a[nIndex] + b[nIndex];
  }

Código C para um núcleo OpenCL:

  __kernel void
  vectorAdd(__global const float * a,
            __global const float * b,
            __global       float * c)
  {
      // Indice do vetor
      int nIndex = get_global_id(0);
      c[nIndex] = a[nIndex] + b[nIndex];
  }

Código C para o HOST CUDA:

Não existe ainda um compilador OpenCL

/* main routine that executes on the host */
int main(void)
{
 float *a_h, *a_d; // Pointer to host & device arrays
 const int N = 10; // Number of elements in arrays
 size_t size = N * sizeof(float);
 a_h = (float *)malloc(size); // Allocate array on host
 cudaMalloc((void **) &a_d, size); // Allocate array on device
 // Initialize host array and copy it to CUDA device
 for (int i=0; i<N; i++) a_h[i] = (float)i;
 cudaMemcpy(a_d, a_h, size, cudaMemcpyHostToDevice);
 // Do calculation on device:
 int block_size = 3;
 int n_blocks = N/block_size + (N % block_size == 0 ? 0:1);
 square_array <<< n_blocks, block_size >>> (a_d, N);
 // Retrieve result from device and store it in host array
 cudaMemcpy(a_h, a_d, sizeof(float)*N, cudaMemcpyDeviceToHost);
 // Print results
 for (int i=0; i<N; i++) printf("%d %f\n", i, a_h[i]);
 // Cleanup
 free(a_h); cudaFree(a_d);
}

Código C para o HOST OpenCL:

# include <stdio.h>
# include <cl.h>
int main(void){ // main routine that executes on the host
 float *a_h; // Pointer to host & device arrays
 const int N = 10; // Number of elements in arrays
 size_t size = N * sizeof(float);
 a_h = (float *)malloc(size); // Allocate array on host
 // Initialize host array and copy it to CUDA device
 for (int i=0; i<N; i++) a_h[i] = (float)i;
 // create the OpenCL context on a GPU device
 cl_context context = clCreateContextFromType(0, // (must be 0)
                      CL_DEVICE_TYPE_ GPU, NULL, // error callback
                      NULL, // user data
                      NULL); // error code
 // get the list of GPU devices associated with context
 size_t cb;
 clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, NULL, &cb);
 cl_device_id *devices = malloc(cb);
 clGetContextInfo(context, CL_CONTEXT_DEVICES, cb, devices, NULL);
 // create a command-queue
 cl_cmd_queue cmd_queue = clCreateCommandQueue(context,
                          devices[0], 0, // default options
                          NULL); // error code
 cl_mem memobjs[1];
 // allocate input buffer memory objects
 memobjs[0] = clCreateBuffer(context,
              CL_MEM_READ_ONLY | // flags
              CL_MEM_COPY_HOST_PTR,
              sizeof(cl_float)*n, // size
              a_h, // host pointer
              NULL); // error code
 // create the program
 cl_program program = clCreateProgramWithSource(
                      context,
                      1, // string count
                      &program_source, // program strings
                      NULL, // string lengths
                      NULL); // error code
 // build the program
 cl_int err = clBuildProgram(program,
              0, // num devices in device list
              NULL, // device list
              NULL, // options
              NULL, // notifier callback function ptr
              NULL); // error code
 // create the kernel
 cl_kernel kernel = clCreateKernel(program, “square_array”, NULL);
 // set “a” vector argument
 err = clSetKernelArg(kernel,
       0, // argument index
       (void *)&memobjs[0], // argument data
       sizeof(cl_mem)); // argument data size
 size_t global_work_size[1] = n; // set work-item dimensions
 // execute kernel
 err = clEnqueueNDRangeKernel(cmd_queue, kernel,
       1, // Work dimensions
       NULL, // must be NULL (work offset)
       global_work_size,
       NULL, // automatic local work size
       0, // no events to wait on
       NULL, // event list
       NULL); // event for this kernel
 // read output array
 err = clEnqueueReadBuffer( context, memobjs[0],
       CL_TRUE, // blocking
       0, // offset
       n*sizeof(cl_float), // size
       a_h, // pointer
       0, NULL, NULL); // events
 // Print results
 for (int i=0; i<N; i++) printf("%d %f\n", i, a_h[i]);
 // Cleanup Host Memory
 free(a_h);
}