OpenCL

O que é

Pela definição da própria Khronos:

“Open Computing Language (OpenCL) é um padrão aberto e isento de royalties para multiplataforma e programação paralela ou seja ele pode ser utilizado em vários tipos de sistemas operacionais, podendo ser usado em supercomputadores, servidores em nuvem, computadores pessoais, dispositivos móveis e plataformas embarcadas. Ele melhora muito a velocidade e a capacidade de resposta de um amplo espectro de aplicativos em várias categorias de mercado, incluindo ferramentas criativas profissionais, softwares científicos e médicos, processamento de visão e treinamento e interfaces de redes neurais.”

Origem

Surgiu em 2008 sendo desenvolvido pela Apple e tendo o Mac-OS-X-10.6 como o primeiro grande SO para apoiar o OpenCL em 2009. Atualmente o OpenCL é mantido pela Khronos Group sendo a Khronos um consórcio de diversas empresas para a produção de padrões abertos.

Funcionamento

O OpenCL fornece uma API para que os programas de software possam realizar processamento paralelo utilizando vários processadores simultaneamente, sendo assim  a eficiência de processamento aumentada pela distribuição da carga entre vários processadores o que causa uma melhora de desempenho de um programa, também é muito usado para acessar a GPU para realizar tarefas gerais de computação.

Características

No OpenCL o runtime fica responsável por criar as instâncias necessárias para o processamento de todo conjunto de dados, não precisando informal ao kernel o tamanho do conjunto de dados.

Arquitetura

Na sua arquitetura o padrão proposto por uma abstração de baixo nível de hardware, sua arquitetura possui 4 modelos, modelo de plataforma, modelo de execução, modelo de programação e modelo de memória.

Modelo de Plataforma

Esse modelo descreve as entidades presentes no ambiente computacional OpenCL, que é integrado por um hospedeiro que agrega um ou mais dispositivos, com esses dispositivos possuindo uma ou mais unidades de computação que são compostas de um ou mais elementos de processamento.

Fonte: César L.B., Luiz G., Gerson G. H.

Modelo de Execução

Esse modelo descreve a instanciação de kernels e a identificação das instâncias. O N-Dimensional Range (NDRange) é um espaço de 1,2 ou 3 dimensões no qual o kernel é executado, as instâncias do kernel são chamados de item de trabalho sendo um índice para cada dimensão do espaço de índices.

Fonte: César L.B., Luiz G., Gerson G. H.

Modelo de Programação

Descreve possíveis métodos de escrita e execução do código opencl e os kernels podem ser executados de modo de paralelismo de dados ou de paralelismo de tarefas, sendo que no paralelismo de dados os itens de trabalho são instanciados para a execução do kernel e no de tarefas um único item de trabalho é instanciado para a execução o que permite vários kernels sobre um conjunto de dados ou muitos conjuntos de dados distintos.

Modelo de Memória

O modelo de memória é dividido em quatro categorias, sendo elas a memória global, a memória local, a memória privada e a memória constante, sendo que a memória global é compartilhada entre todos os itens de trabalho e também é permitido escrita e leitura, a memória local é compartilhada apenas entre os  itens de um mesmo trabalho e tendo permissão de escrita e leitura, a memória privada é restrita a cada item de trabalho e a memória constante é compartilhada entre todos os itens de trabalho mas apenas com permissão de leitura.

Linguagem OpenCL C

É baseada na especificação C99 da linguagem C.

Tipos de Dados

Existem 2 categorias de dados no OpenCL, as do tipo escalares e tipos vetoriais. Os tipos escalares armazenam informações numéricas em diversos intervalos e com diferentes níveis de precisão. Os tipos vetoriais armazenam múltiplos dados escalares, struct, union, arrays, funções e ponteiros. Os dados são declarados com o tipo do dado seguido da quantidade de componentes no vetor, um exemplo seria float4 que representa um vetor com 4 componentes do tipo float.

Os tipos de dados da linguagem OpenCL C são:

uchar / char

Onde u é o prefixo usado ao invés de unsigned para tipos vetoriais com componentes sem sinal. Os tipos bool, half, size_t e void não suportam vetores.

Qualificadores

No OpenCL C os operadores oferecidos trabalham com dados do tipo escalar e do tipo vetorial. Os operadores são:

Algumas observações importantes, 1º é que uma operação de um vetor e um escalar converte o escalar para vetor e aplica a operação sobre cada componente do vetor. 2º Operações entre vetores requerem que eles sejam do mesmo tamanho. 3º Operações lógicas e relacionais com pelo menos um operando vetorial retornam um vetor do mesmo tipo desde operando. 4º Operadores aritméticos(%, ++, --), operadores lógicos e bit-a-bit não suportam pontos-flutuantes.

Restrições

A OpenCL C estende a linguagem C99 com tipos novos e alguns outros elementos, da mesma forma adiciona restrições sendo alguns deles sobre: não ser permitido ponteiros para função; que argumentos para kernels não podem ser ponteiros para ponteiros; macros e funções com número variável de argumentos não são suportados; os qualificadores extern, static e auto não são suportados; não existe suporte para recursão; ponteiros e arrays com tamanho inferior a 32 bits não são permitidos e tipos vetoriais char2 e uchar 2 também não são permitidos; elementos de uma struct ou union devem pertencer ao mesmo espaço de endereçamento.

Exemplos

Para exemplo, o programa a seguir simplesmente retorna a mensagem “Hello, World” que será gerada no kernel.

__kernel void hello(__global char* string)
{
    string[0] = 'H';
    string[1] = 'e';
    string[2] = 'l';
    string[3] = 'l';
    string[4] = 'o';
    string[5] = ',';
    string[6] = ' ';
    string[7] = 'W';
    string[8] = 'o';
    string[9] = 'r';
    string[10] = 'l';
    string[11] = 'd';
    string[12] = '!';
    string[13] = '\0';
}

No caso o kernel é apenas uma cadeia de caracteres que será enviada ao programa principal.

#include <stdio.h>
#include <stdlib.h>
 
#define CL_USE_DEPRECATED_OPENCL_1_2_APIS
 
#include <CL/cl.h>
 
#define MEM_SIZE (128)
#define MAX_SOURCE_SIZE (0x100000)
 
int main()
{
    cl_device_id device_id = NULL;
    cl_context context = NULL;
    cl_command_queue command_queue = NULL;
    cl_mem memobj = NULL;
    cl_program program = NULL;
    cl_kernel kernel = NULL;
    cl_platform_id platform_id = NULL;
    cl_uint ret_num_devices;
    cl_uint ret_num_platforms;
    cl_int ret;
   
    char string[MEM_SIZE];
		 FILE *fp;
    char fileName[] = "./kernels/hello.cl";
    char *source_str;
    size_t source_size;
   
    /* Load the source code containing the kernel*/
    fp = fopen(fileName, "r");
        if (!fp) {
        fprintf(stderr, "Failed to load kernel.\n");
        exit(1);
    }
    source_str = (char*)malloc(MAX_SOURCE_SIZE);
    source_size = fread(source_str, 1, MAX_SOURCE_SIZE, fp);
    fclose(fp);
   
    /* Get Platform and Device Info */
    ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms);
    ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, &ret_num_devices);
   
    /* Create OpenCL context */
    context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret);
   
    /* Create Command Queue */
    command_queue = clCreateCommandQueue(context, device_id, 0, &ret);
   
    /* Create Memory Buffer */
    memobj = clCreateBuffer(context, CL_MEM_READ_WRITE,MEM_SIZE * sizeof(char), NULL, &ret);
   
    /* Create Kernel Program from the source */
    program = clCreateProgramWithSource(context, 1, (const char **)&source_str,
    (const size_t *)&source_size, &ret);
   
    /* Build Kernel Program */
    ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);
   
    /* Create OpenCL Kernel */
    kernel = clCreateKernel(program, "hello", &ret);
   
    /* Set OpenCL Kernel Parameters */
    ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&memobj);
   
    /* Execute OpenCL Kernel */
    ret = clEnqueueTask(command_queue, kernel, 0, NULL,NULL);
   
    /* Copy results from the memory buffer */
    ret = clEnqueueReadBuffer(command_queue, memobj, CL_TRUE, 0,
    MEM_SIZE * sizeof(char),string, 0, NULL, NULL);
		/* Display Result */
    puts(string);
   
    /* Finalization */
    ret = clFlush(command_queue);
    ret = clFinish(command_queue);
    ret = clReleaseKernel(kernel);
    ret = clReleaseProgram(program);
    ret = clReleaseMemObject(memobj);
    ret = clReleaseCommandQueue(command_queue);
    ret = clReleaseContext(context);
   
    free(source_str);
   
    return 0;
}

Este trecho do código é o responsável por ler e exibir o resultado:

/* Copy results from the memory buffer */
    ret = clEnqueueReadBuffer(command_queue, memobj, CL_TRUE, 0,
    MEM_SIZE * sizeof(char),string, 0, NULL, NULL);
   
    /* Display Result */
    puts(string);

E vale lembrar que dependendo da versão utilizada pode não ser preciso desalocar os recursos, sendo o trecho final do código o responsável por isto:

/* Finalization */
    ret = clFlush(command_queue);
    ret = clFinish(command_queue);
    ret = clReleaseKernel(kernel);
    ret = clReleaseProgram(program);
    ret = clReleaseMemObject(memobj);
    ret = clReleaseCommandQueue(command_queue);
    ret = clReleaseContext(context);
   
    free(source_str);

Materiais

Programação em OpenCL: Uma introdução prática

OpenCL in Action

Introdução ao OpenCL

Acadêmicos: André Guilherme dos Santos e Kelvin Juliano dos Santos