Comparando CPU vs GPU usando OpenCL

Atualmente temos uma demanda muito grande por mais poder de processamento e velocidade em dispositivos tecnológicos. Para acompanhar o ritmo desses avanços, empresas encontram diversos meios para melhorar poder de processamento em dispositivos. Um meio que a Apple Inc. encontrou foi criar a Open Computing Language (OpenCL). No dia 16 de Junho de 2008 a Apple propôs para o Khronos Group para trabalharem no OpenCL. Depois de 5 meses tomando muito café e programando, no dia 8 de Dezembro de 2008, OpenCL 1.0 foi divulgado para a comunidade.

OpenCL é uma API de baixo nível para programação paralela de diversos tipos de processadores que podem ser encontrados em computadores pessoais, servidores, dispositivos mobile, como também em sistemas embarcados. A linguagem de programação usada pelo OpenCL é uma linguagem parecida com C e pode ser usada em plataformas heterogêneas que contêm CPUs, GPUs, e processadores de fabricantes como NXP, Nvidia, Intel, AMD e IBM. O propósito do OpenCL é acelerar e melhorar a capacidade de resposta das mais diversas aplicações encontradas no mercado, desde jogos e entretenimento até aplicações médicas e científicas.

Neste artigo iremos experimentar o OpenCL usando o SOM Apalis iMX6Q da Toradex, comparando duas aplicações. Uma delas irá rodar na GPU do processador e a outra na própria CPU. No final compartilharemos os resultados obtidos.

Hardware Usado

O computador em módulo da Toradex Apalis iMX6Q é baseado no processador iMX6Quad da NXP, o qual oferece recursos de processamento eficientes particularmente adequados a aplicações gráficas e multimídia. O processador tem quatro núcleos ARM® Cortex-A9® com até 800MHz por núcleo. Além do processador, o computador em módulo conta com memórias de 2GB DDR3 RAM (64bit) e 4GB eMMC Flash.

Os Melhores Treinamentos sobre Sistemas embarcados e IoT

Cursos com professores qualificados para acelerar sua carreira e projetos

MATRÍCULE-SE AGORA

Focando em propósitos gráficos e de multimídia, o processador também oferece uma GPU 3D Vivante GC2000 que é capaz de suportar OpenCL EP (Embedded Profile) versão 1.1, portanto, podemos usar o poder de processamento da GPU do Apalis iMX6Q em diversas aplicações.

Suporte à  OpenCL na imagem de Linux embarcado da Toradex 

Partimos do ponto onde já possuímos um ambiente de geração de imagens OpenEmbedded já configurado e pronto para gerar uma imagem para Apalis iMX6Q. Isso pode ser realizado seguindo o artigo no Portal de Desenvolvedores da Toradex.

Para gerar uma imagem de Linux embarcado que suporta OpenCL EP 1.1 e também inclui suas bibliotecas, é necessário realizar alguns passos adicionais descritos adiante.

Primeiro, edite o arquivo do seguinte diretório:

~/meta-toradex/recipes-fsl/packagegroups/pakcagegroup-fsl-tools-gpu.bbappend 

Adicionando o seguinte conteúdo:

SOC_TOOLS_GPU_append_mx6 = " \ 
	libopencl-mx6 \ 
	libgles-mx6 \ 
" 

Também adicione o pacote imx-gpu-viv no arquivo local.conf:

IMAGE_INSTALL_append = "imx-gpu-viv" 

E inicie o processo de geração de uma imagem Desktop:

bitbake angstrom-lxde-image 

Código da GPU e CPU

Todo o código neste artigo pode ser encontrado no GitHub. 

Como exemplo, usamos duas aplicações que basicamente somam vetores. O primeiro código é executado na GPU e o segundo na CPU. O tempo consumido é mostrado no terminal quando as aplicações se encerram. O header necessário para usar OpenCL é cl.h e pode ser encontrado em /usr/include/CL no rootfs. As bibliotecas necessárias para rodar os programas são libGAL.so, que faz as chamadas para a GPU Vivante GC2000, e libOpenCL.so, que corresponde às implementações e interfaces do OpenCL. Ambas podem ser encontradas em /usr/lib. 

Para o cálculo do tempo consumido, criamos uma queue, com profilling habilitado, e então lemos as informações de profilling no final do programa.

Segue o código do OpenCL:

//************************************************************ 
// Demo OpenCL application to compute a simple vector addition 
// computation between 2 arrays on the GPU 
// ************************************************************ 
#include <stdio.h> 
#include <stdlib.h> 
#include <time.h> 
#include <CL/cl.h> 
// 
// OpenCL source code 
const char* OpenCLSource[] = { 
  "__kernel void VectorAdd(__global int* c, __global int* a,__global int* b)", 
  "{", 
  " // Index of the elements to add \n", 
  " unsigned int n = get_global_id(0);", 
  " // Sum the nth element of vectors a and b and store in c \n", 
  " c[n] = a[n] + b[n];", 
  "}" 
};
 
// Some interesting data for the vectors 
Int InitialData1[80] = {37,50,54,50,56,0,43,43,74,71,32,36,16,43,56,100,50,25,15,17,37,50,54,50,56,0,43,43,74,71,32,36,16,43,56,100,50,25,15,17,37,50,54,50,56,0,43,43,74,71,32,36,16,43,56,100,50,25,15,17,37,50,54,50,56,0,43,43,74,71,32,36,16,43,56,100,50,25,15,17}; 
int InitialData2[80] = {35,51,54,58,55,32,36,69,27,39,35,40,16,44,55,14,58,75,18,15,35,51,54,58,55,32,36,69,27,39,35,40,16,44,55,14,58,75,18,15,35,51,54,58,55,32,36,69,27,39,35,40,16,44,55,14,58,75,18,15,35,51,54,58,55,32,36,69,27,39,35,40,16,44,55,14,58,75,18,15}; 
// Number of elements in the vectors to be added 
#define SIZE 600000 

// Main function 
// ************************************************************ 
int main(int argc, char **argv) 
{  
     // Two integer source vectors in Host memory 
     int HostVector1[SIZE], HostVector2[SIZE]; 
     //Output Vector 
     int HostOutputVector[SIZE]; 
     // Initialize with some interesting repeating data 
     for(int c = 0; c < SIZE; c++) 
     { 
          HostVector1[c] = InitialData1[c%20]; 
          HostVector2[c] = InitialData2[c%20]; 
          HostOutputVector[c] = 0; 
     } 
     //Get an OpenCL platform 
     cl_platform_id cpPlatform; 
     clGetPlatformIDs(1, &cpPlatform, NULL); 
     // Get a GPU device 
     cl_device_id cdDevice; 
     clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_GPU, 1, &cdDevice, NULL); 
     char cBuffer[1024]; 
     clGetDeviceInfo(cdDevice, CL_DEVICE_NAME, sizeof(cBuffer), &cBuffer, NULL); 
     printf("CL_DEVICE_NAME: %s\n", cBuffer); 
     clGetDeviceInfo(cdDevice, CL_DRIVER_VERSION, sizeof(cBuffer), &cBuffer, NULL); 
     printf("CL_DRIVER_VERSION: %s\n\n", cBuffer); 
     // Create a context to run OpenCL enabled GPU 
     cl_context GPUContext = clCreateContextFromType(0, CL_DEVICE_TYPE_GPU, NULL, NULL, NULL);      
     // Create a command-queue on the GPU device 
     cl_command_queue cqCommandQueue = clCreateCommandQueue(GPUContext, cdDevice, CL_QUEUE_PROFILING_ENABLE, NULL); 
     // Allocate GPU memory for source vectors AND initialize from CPU memory 
     cl_mem GPUVector1 = clCreateBuffer(GPUContext, CL_MEM_READ_ONLY | 
     CL_MEM_COPY_HOST_PTR, sizeof(int) * SIZE, HostVector1, NULL); 
     cl_mem GPUVector2 = clCreateBuffer(GPUContext, CL_MEM_READ_ONLY | 
     CL_MEM_COPY_HOST_PTR, sizeof(int) * SIZE, HostVector2, NULL); 
     // Allocate output memory on GPU 
     cl_mem GPUOutputVector = clCreateBuffer(GPUContext, CL_MEM_WRITE_ONLY, 
     sizeof(int) * SIZE, NULL, NULL); 
     // Create OpenCL program with source code 
     cl_program OpenCLProgram = clCreateProgramWithSource(GPUContext, 7, OpenCLSource, NULL, NULL); 
     // Build the program (OpenCL JIT compilation) 
     clBuildProgram(OpenCLProgram, 0, NULL, NULL, NULL, NULL); 
     // Create a handle to the compiled OpenCL function (Kernel) 
     cl_kernel OpenCLVectorAdd = clCreateKernel(OpenCLProgram, "VectorAdd", NULL); 
     // In the next step we associate the GPU memory with the Kernel arguments 
     clSetKernelArg(OpenCLVectorAdd, 0, sizeof(cl_mem), (void*)&GPUOutputVector); 
     clSetKernelArg(OpenCLVectorAdd, 1, sizeof(cl_mem), (void*)&GPUVector1); 
     clSetKernelArg(OpenCLVectorAdd, 2, sizeof(cl_mem), (void*)&GPUVector2); 
      
     //create event 
     cl_event event = clCreateUserEvent(GPUContext, NULL); 
      
     // Launch the Kernel on the GPU 
     // This kernel only uses global data 
     size_t WorkSize[1] = {SIZE}; // one dimensional Range 
     clEnqueueNDRangeKernel(cqCommandQueue, OpenCLVectorAdd, 1, NULL, WorkSize, NULL, 0, NULL, &event); 
     // Copy the output in GPU memory back to CPU memory 
     clEnqueueReadBuffer(cqCommandQueue, GPUOutputVector, CL_TRUE, 0, 
     SIZE * sizeof(int), HostOutputVector, 0, NULL, NULL); 
     // Cleanup 
     clReleaseKernel(OpenCLVectorAdd); 
     clReleaseProgram(OpenCLProgram); 
     clReleaseCommandQueue(cqCommandQueue); 
     clReleaseContext(GPUContext); 
     clReleaseMemObject(GPUVector1); 
     clReleaseMemObject(GPUVector2); 
     clReleaseMemObject(GPUOutputVector);     
      
     clWaitForEvents(1, &event); 
     cl_ulong start = 0, end = 0; 
     double total_time;      
      
     clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &start, NULL); 
     clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &end, NULL); 
      
     total_time = end - start;      
              
     printf("\nExecution time in milliseconds = %0.3f ms", (total_time / 1000000.0) ); 
     printf("\nExecution time in seconds = %0.3f s\n\n", ((total_time / 1000000.0))/1000 );           
           
     return 0; 
} 

O código da CPU por sua vez é um código escrito puramente em C que realiza a mesma soma de vetores do programa anterior. Para calcular o tempo consumido usamos a biblioteca time.h. O código é visto a seguir:

#include <stdio.h> 
#include <stdlib.h> 
#include <time.h>  
 
int InitialData1[80] = {37,50,54,50,56,0,43,43,74,71,32,36,16,43,56,100,50,25,15,17,37,50,54,50,56,0,43,43,74,71,32,36,16,43,56,100,50,25,15,17,37,50,54,50,56,0,43,43,74,71,32,36,16,43,56,100,50,25,15,17,37,50,54,50,56,0,43,43,74,71,32,36,16,43,56,100,50,25,15,17}; 
int InitialData2[80] = {35,51,54,58,55,32,36,69,27,39,35,40,16,44,55,14,58,75,18,15,35,51,54,58,55,32,36,69,27,39,35,40,16,44,55,14,58,75,18,15,35,51,54,58,55,32,36,69,27,39,35,40,16,44,55,14,58,75,18,15,35,51,54,58,55,32,36,69,27,39,35,40,16,44,55,14,58,75,18,15}; 
 
#define SIZE 600000 
 
int main(int argc, char **argv) 
{ 
  time_t start, stop; 
  clock_t ticks; 
     
  time(&start);     
  // Two integer source vectors in Host memory 
  int HostVector1[SIZE], HostVector2[SIZE]; 
  //Output Vector 
  int HostOutputVector[SIZE]; 
  // Initialize with some interesting repeating data 
  //int n; 
  for(int c = 0; c < SIZE; c++) 
  { 
    HostVector1[c] = InitialData1[c%20]; 
    HostVector2[c] = InitialData2[c%20]; 
    HostOutputVector[c] = 0; 
  } 
      
  for(int i = 0; i < SIZE; i++) 
  { 
    HostOutputVector[i] = HostVector1[i] + HostVector2[i]; 
    ticks = clock(); 
  }      
 
  time(&stop); 
 
  printf("\nExecution time in miliseconds = %0.3f ms",((double)ticks/CLOCKS_PER_SEC)*1000); 
 
  printf("\nExecution time in seconds = %0.3f s\n\n", (double)ticks/CLOCKS_PER_SEC); 
 
  return 0; 
} 

Cross Compilando as Aplicações

Um mesmo Makefile pode ser usado para realizar compilacão-cruzada de ambas as aplicações: CPU e GPU. Atente para as três variáveis seguintes do Makefile que precisam ser editadas de acordo com seu sistema:

• ROOTFS_DIR – Diretório do sysroots do Apalis iMX6;
• APPNAME – Nome da aplicação;
• TOOLCHAIN – Diretório da toolchain para cross-compilar.

export ARCH=arm 
export ROOTFS_DIR=/usr/local/toradex-linux-v2.5/oe-core/build/out-glibc/sysroots/apalis-imx6 
 
APPNAME = proc_sample 
TOOLCHAIN = /home/prjs/toolchain/gcc-linaro 
 
CROSS_COMPILER = $(TOOLCHAIN)/bin/arm-linux-gnueabihf- 
CC= $(CROSS_COMPILER)gcc 
DEL_FILE = rm -rf 
CP_FILE = cp -rf 
TARGET_PATH_LIB = $(ROOTFS_DIR)/usr/lib 
TARGET_PATH_INCLUDE = $(ROOTFS_DIR)/usr/include 
CFLAGS = -DLINUX -DUSE_SOC_MX6 -Wall -std=c99 -O2 -fsigned-char -march=armv7-a -mfpu=neon -DEGL_API_FB -DGPU_TYPE_VIV -DGL_GLEXT_PROTOTYPES -DENABLE_GPU_RENDER_20 -I../include -I$(TARGET_PATH_INCLUDE) 
LFLAGS = -Wl,--library-path=$(TARGET_PATH_LIB),-rpath-link=$(TARGET_PATH_LIB) -lm -lglib-2.0 -lOpenCL -lCLC -ldl -lpthread 
OBJECTS = $(APPNAME).o 

first: all 

all: $(APPNAME) 
	$(APPNAME): $(OBJECTS) 
	$(CC) $(LFLAGS) -o $(APPNAME) $(OBJECTS) 
	$(APPNAME).o: $(APPNAME).c 
	$(CC) $(CFLAGS) -c -o $(APPNAME).o $(APPNAME).c 

clean: 
	$(DEL_FILE) $(APPNAME)

Salve o Makefile no mesmo diretório da aplicação e execute make. Copie os binários para o módulo Apalis iMX6 da forma que julgar apropriado (SCP, FTP, etc).

Resultados Finais

Depois de executar ambas as aplicações obtivemos os seguintes resultados:

### Processor time 
Execution time in miliseconds = 778.999 ms 
Execution time in seconds = 0.779 s  
 
### GPU time  
Execution time in milliseconds = 12.324 ms 
Execution time in seconds = 0.012 s 

Baseando-se nos resultados podemos ver claramente que conseguimos acelerar o cálculo de soma de vetores usando o poder de processamento da GPU do módulo Apalis iMX6Q por quase 65 vezes (64.92, para ser mais exato)!

Conclusão

Aqueles que desejam aproveitar-se da GPU Vivante GC2000 presente no Apalis iMX6 podem, além de outros métodos, usar OpenCL para aumentar o poder de processamento em rotinas que exigem muitos cálculos vetoriais e matriciais. Com os recursos do OpenCL é possível rodar aplicações em dispositivos desde placas de vídeo e supercomputadores até sistemas embarcados como visto neste artigo. Poderiam até mesmo ir além, por exemplo, utilizando OpenCL com OpenCV para aumentar a performance de aplicações de visão computacional. Este artigo serve de exemplo para infindas possibilidades de aplicações que uma empresa pode desenvolver.

Referências