ÍNDICE DE CONTEÚDO
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.
Embarcados Experience 2024: Evento Presencial
Participe do Embarcados Experience 2024 em São Paulo. Conhecimento técnico, palestras, workshops e oportunidade de networking com profissionais experientes.
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:
|
1 |
~/meta-toradex/recipes-fsl/packagegroups/pakcagegroup-fsl-tools-gpu.bbappend |
Adicionando o seguinte conteúdo:
|
1 2 3 4 |
SOC_TOOLS_GPU_append_mx6 = " \ libopencl-mx6 \ libgles-mx6 \ " |
Também adicione o pacote imx-gpu-viv no arquivo local.conf:
|
1 |
IMAGE_INSTALL_append = "imx-gpu-viv" |
E inicie o processo de geração de uma imagem Desktop:
|
1 |
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:
|
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 82 83 84 85 86 87 88 89 90 91 92 93 94 95 96 97 98 99 100 101 102 103 104 105 106 107 108 |
//************************************************************ // 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:
|
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 |
#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.
|
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 |
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:
|
1 2 3 4 5 6 7 |
### 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
Olá, no loop da CPU tem uma chamada ticks = clock(), mas ela deveria estar fora do loop, após ele ter terminado. Você poderia corrigir ? Obrigado,
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… Leia mais »
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