Ver código fonte para OpenCL

{{não confundir com|OpenGL}}
[[Ficheiro:Clinfo screenshot.png|miniaturadaimagem|Imagem ilustrativa]]
'''OpenCL''' ('''Open''' '''C'''omputing '''L'''anguage) é uma arquitetura para escrever programas que funcionam em plataformas heterogêneas, consistindo em [[Central processing unit|CPUs]], [[Graphics processing unit|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.

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 [[Graphics processing unit|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<ref>{{citar comunicado de imprensa|url=http://www.khronos.org/news/press/releases/khronos_launches_heterogeneous_computing_initiative/ |título=Khronos Launches Heterogeneous Computing Initiative |publicado=Khronos Group  |data=2008-06-16}}</ref> 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.<ref name=macWorld>
{{citar web
|url=http://www.macworld.com/article/136921/2008/11/opencl.html?lsrc=top_2
|título=OpenCL gets touted in Texas
|publicado=MacWorld
|data=2008-11-20
}}
</ref>

Esta especificação técnica foi revisada por membros da Khronos e aprovada como release pública em 8 de dezembro de 2008.<ref name="khronosGroup">
{{citar comunicado de imprensa
|url=http://www.khronos.org/news/press/releases/the_khronos_group_releases_opencl_1.0_specification/
|título=The Khronos Group Releases OpenCL 1.0 Specification
|publicado=Khronos Group
|data=2008-12-08
}}
</ref>

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:<ref name=pressrelease>{{citar comunicado de imprensa|url=http://www.apple.com/pr/library/2008/06/09snowleopard.html  |título=Apple Previews Mac OS X Snow Leopard to Developers |publicado=Apple Inc. |data=2008-06-09}}</ref>
{{Quote|
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 [[AMD FireStream#Software Development Kit|Stream framework]].<ref name=AMDpressrelease>
{{citar comunicado de imprensa
|url=http://www.amd.com/us-en/Corporate/VirtualPressRoom/0,,51_104_543~127451,00.html
|título=AMD Drives Adoption of Industry Standards in GPGPU Software Development
|publicado=AMD}}
</ref><ref name=eweekAMD>
{{citar web
|url=http://www.eweek.com/c/a/Desktops-and-Notebooks/AMD-Backing-OpenCL-and-Microsoft-DirectX-11/
|título=AMD Backs OpenCL, Microsoft DirectX 11
|publicado=eWeek}}
</ref>

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.<ref name="RapidMindHPCWire">
{{citar web
|url=http://www.hpcwire.com/topic/applications/RapidMind_Embraces_Open_Source_and_Standards_Projects.html
|título=HPCWire: RapidMind Embraces Open Source and Standards Projects
|publicado=HPCWire
|acessodata=2009-04-10
|arquivourl=https://web.archive.org/web/20081218113648/http://www.hpcwire.com/topic/applications/RapidMind_Embraces_Open_Source_and_Standards_Projects.html
|arquivodata=2008-12-18
|urlmorta=yes
}}
</ref>

[[Nvidia]] anunciou em 9 de dezembro de 2008 agregar suporte completo à especificação OpenCL 1.0 ao seu GPU Computing Toolkit.<ref name="Nvidia Press Release 2008-12-09">
{{citar comunicado de imprensa
|url=http://www.nvidia.com/object/io_1228825271885.html
|título=NVIDIA Adds OpenCL To Its Industry Leading GPU Computing Toolkit
|publicado=Nvidia}}</ref>

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:

<syntaxhighlight lang="c">
  __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];
  }
</syntaxhighlight>

Código C para um núcleo OpenCL:
<syntaxhighlight lang="c">
  __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];
  }
</syntaxhighlight>

Código<ref name=gpumodeling>{{citar web|url=http://gpumodeling.blogspot.com/2009/04/opencl-and-cuda.html  |título=GPU Modeling |publicado=Blogspot |data=2009-04-21}}</ref>  C para o HOST CUDA:

''Não existe ainda um compilador OpenCL''

<syntaxhighlight lang="c">

/* 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);
}
</syntaxhighlight>

Código<ref name="gpumodeling" /> C para o HOST OpenCL:

<syntaxhighlight lang="c">
# 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);
}
</syntaxhighlight>

{{esboço|programação}}
{{referências}}

{{Padronizações da Khronos Group}}

[[Categoria:Linguagens de programação]]
[[Categoria:Padrões gráficos]]