Comparando CPU vs GPU usando OpenCL

CPU GPU 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.

 

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

 

https://www.khronos.org/opencl/

https://en.wikipedia.org/wiki/OpenCL

http://www.drdobbs.com/parallel/a-gentle-introduction-to-opencl/231002854

https://community.freescale.com/docs/DOC-93984

https://community.freescale.com/docs/DOC-100694

http://developer.toradex.com/products/apalis-imx6

https://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/clGetEventProfilingInfo.html

http://parallelis.com/how-to-measure-opencl-kernel-execution-time/

https://software.intel.com/en-us/articles/intel-sdk-for-opencl-applications-performance-debugging-intro

NEWSLETTER

Receba os melhores conteúdos sobre sistemas eletrônicos embarcados, dicas, tutoriais e promoções.

Obrigado! Sua inscrição foi um sucesso.

Ops, algo deu errado. Por favor tente novamente.

Licença Creative Commons Esta obra está licenciada com uma Licença Creative Commons Atribuição-CompartilhaIgual 4.0 Internacional.

Giovanni Bauermeister
É formado em Técnico em Mecatrônica pelo SENAI Roberto Mange de Campinas, onde teve seus primeiros contatos com microcontroladores, eletrônica e programação. Em 2016 obteve a graduação em Engenharia de Controle e Automação pela UNISAL, também em Campinas. Durante a graduação, teve a oportunidade de ingressar como estagiário na empresa Toradex, onde atuou e desenvolveu habilidades em sistemas Linux embarcado. Participou do movimento maker, contribuindo com tutoriais e suporte ao cliente na Filipeflop. Atualmente atua na área de desenvolvimento de projetos.

2
Deixe um comentário

avatar
 
1 Comment threads
1 Thread replies
0 Followers
 
Most reacted comment
Hottest comment thread
2 Comment authors
Giovanni BauermeisterVinicius De Araujo Barboza Recent comment authors
  Notificações  
recentes antigos mais votados
Notificar
Vinicius De Araujo Barboza
Visitante
Vinicius de Araujo Barboza

Giovanni!

Usei esse mesmo teste como comparativo entre CPU e GPU de um Samsung ARM Chromebook alterando algumas coisas no Makefile e obtive valores consistentes com os seus.

Achei que seria interessante podermos comparar o HostOutputVector do código executado na CPU e GPU, mas quando tento o print desse vetor no código da GPU tenho apenas zeros (0 0 0 .. 0 0 0) - já o código da CPU está legal.

Consegue me ajudar? Talvez eu tenha esquecido algo na hora de carregar o buffer da GPU no host, afinal só adicionei umas 2 linhas para um for-loop e printf.

Valeu!

Giovanni Bauermeister
Visitante
Giovanni Bauermeister

Olá Vinicius! Que bom que conseguiu executar o artigo e chegar nos resultados! Para escrever o código do OpenCL tomei como base o seguinte tutorial: https://community.nxp.com/docs/DOC-93984 Originalmente o código "printava" os vetores mas não mostrava o tempo das operações. Eu tirei essa parte do código e modifiquei também pra fazer o cálculo do tempo. Mas se você olhar no original logo após clReleaseMemObject(GPUOutputVector); tem o for com o printf clReleaseMemObject(GPUOutputVector); for( int i =0 ; i < SIZE; i++) printf("[%d + %d = %d]n",HostVector1[i], HostVector2[i], HostOutputVector[i]); return 0; talvez isso te ajude