Experimenting with OpenCL on Apalis iMX6Q

Friday, December 2, 2016
Apalis
Introduction
Nowadays, technological devices are demanding increasingly higher processing power and speed than ever before. To cope with the advances in technology, companies create various ways to achieve better processing performance. One such way that Apple Inc. found was to create the Open Computing Language (OpenCL). On June 16, 2008, Apple submitted a proposal to the Khronos Group to work on OpenCL. After five months of work, OpenCL 1.0 was released to the public on December 8, 2008.
OpenCL is a low level API for parallel programming of diverse processors found in personal computers, servers, mobile devices, as well as embedded devices. The programming language for OpenCL is a C-like language. It is made to work in heterogeneous platforms containing CPU's, GPU's, and processors from such popular manufacturers as NXP®, NVIDIA®, Intel®, AMD, IBM, etc. The purpose of OpenCL is to improve the speed and responsiveness of a wide range of applications such as gaming and, entertainment, as well as scientific and medical software.
In this post, we experiment with OpenCL using Toradex's Apalis iMX6Q SoM to compare two applications - one of them running on the GPU and the other on the CPU. At the end we share the results found in this experiment.
OpenCL on Apalis iMX6Q
Hardware used
Toradex's Apalis iMX6Q SoM is based on NXP's iMX6Quad processor which offers efficient processing capabilities particularly suited to multimedia applications. The processor has four ARM® Cortex®-A9 cores up to 800 MHz per core. In addition to the processor, the Apalis SoM also offers up to 2GB DDR3 RAM (64 Bit) and 4GB eMMC FLASH.
Focusing on graphics and multimedia purposes, this processor also offers a Vivante GC2000 3D GPU which is capable of running OpenCL EP (Embedded Profile); therefore, we can use the i.MX6Q GPU processing power in any program.
OpenCL support in Toradex Embedded Linux image
We start from the point where we already have an OpenEmbedded build system configured ready to build an image for Apalis iMX6. This can be achieved following our OpenEmbedded (core) article.
For building an Embedded Linux image which supports OpenCL and includes its libraries, some additional steps need to be taken.
First, edit the following file in the directory;
/meta-toradex/recipes-fsl/packagegroups/packagegroup-fsl-tools-gpu.bbappend
adding the following content:
SOC_TOOLS_GPU_append_mx6 = " \
    libopencl-mx6 \
    libgles-mx6 \
"

Also, add the package imx-gpu-viv in local.conf file:
IMAGE_INSTALL_append = "imx-gpu-viv"

Now, build a Desktop based image:
bitbake angstrom-lxde-image
GPU and CPU code
All the code in this post can be found on GitHub.
As an example, we used two basic applications which perform a simple sum of arrays. The first code runs on the GPU and the second on the CPU. The consumed time is printed after the applications finish. The header needed to use OpenCL is cl.h which can be found at /usr/include/CL in the roofts. The libraries needed to link the program are libGAL.so and libOpenCL.so. They can be found at /usr/lib.
To calculate the consumed time, we created a queue with profiling enabled and got the profiling data at the end.
Follow the OpenCL code:
//************************************************************
// Demo OpenCL application to compute a simple vector addition
// computation between 2 arrays on the GPU
// ************************************************************
#include 
#include 
#include 
#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, &amp;cpPlatform, NULL);
     // Get a GPU device
     cl_device_id cdDevice;
     clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_GPU, 1, &amp;cdDevice, NULL);
     char cBuffer[1024];
     clGetDeviceInfo(cdDevice, CL_DEVICE_NAME, sizeof(cBuffer), &amp;cBuffer, NULL);
     printf("CL_DEVICE_NAME: %s\n", cBuffer);
     clGetDeviceInfo(cdDevice, CL_DRIVER_VERSION, sizeof(cBuffer), &amp;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*)&amp;GPUOutputVector);
     clSetKernelArg(OpenCLVectorAdd, 1, sizeof(cl_mem), (void*)&amp;GPUVector1);
     clSetKernelArg(OpenCLVectorAdd, 2, sizeof(cl_mem), (void*)&amp;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, &amp;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, &amp;event);
     cl_ulong start = 0, end = 0;
     double total_time;     
 
     clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &amp;start, NULL);
     clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &amp;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;
}

The CPU code is a simple C program which computes the same sum of arrays from above. To calculate the time consumed, we used the library time.h. The code is as follows:
#include 
#include 
#include  
 
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(&amp;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(&amp;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 compiling the applications
One Makefile can be used to cross compile both GPU and CPU applications. The following three variables need your attention. You can change them according to your system:
  • ROOTFS_DIR -> the sysroots path of Apalis iMX6
  • APPNAME -> the name of you application
  • TOOLCHAIN -> the the cross compiler toolchain path
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)

Save the Makefile in the same folder of your applications and run make. 
Copy the generated binaries to Apalis iMX6.
End results
After running both applications we got the following results:
### 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

Based on these results, we can clearly see that we were able to speed up the array sum by using OpenCL with Apalis iMX6Q GPU's processing capabilities.
Conclusion
Customers willing to take advantage of Apalis iMX6Q GPU can, among other methods, use OpenCL to increase computing power. With OpenCL capabilities, it is possible to run code in devices ranging from graphics cards to supercomputers as well as embedded devices, as seen in this post. Customers could also go further combining, for example, OpenCL with OpenCV to increase performance in computer vision applications. This can be used as an example of endless possibilities of applications that a company can develop.
References

This blog post was originally featured on Embarcados.com in Portuguese. See here.
Author: Giovanni Bauermeister, Toradex Brasil

Comments

Post a Comment

Popular posts from this blog

Toradex WinCE Libraries' Evolution: The New Library Approach Explained

Image
Tuesday, November 8, 2016 Further information about our libraries including demo application downloads can be found in our  developer center . Overview Over 10 years ago, we wrote some small pieces of WinCE code in order to demonstrate different features of our  computer on modules  (CoM), or  system on modules  (SoM); for example: GPIOs, I2C, Analog Inputs (ADC), etc. With the pure demo purpose in our minds we didn’t pay attention to things like maintenance, compatibility across different modules, etc. However, these demos were used by a huge Windows Embedded Compact customer base. We received more and more support and feature requests which we then started to integrate into the demo code. As a result, we got the  library package  for our  PXA  and  Tegra ®  based SoMs  as it exists today. These libraries are a reliable and easy way of using our modules. But it wasn’t an easy way to get there with the many maintenance iterations and bending the API to make use of newer featu
Read more

Patching Kernel in OpenEmbedded

Image
Thursday, January 19, 2017 Introduction Despite  OpenEmbedded / Yocto  being nice tools to generate your image, using the build system to compile and test tiny alterations applied to independent packages, although possible, can become confusing and slow. My suggestion to alter the kernel is to compile it separately, test all the changes made by the end of the project, and add the changes to OpenEmbedded/Yocto. To illustrate the process, we will add support to the inertial measurement unit  mma8653  accelerometer by  NXP  (formerly Freescale) to the kernel 4.4.14 used by the Toradex computer on module  Colibri VF61 . Procedure The first thing to do is to have OpenEmbedded and its prerequisites installed. You can follow  this article 's  prerequisites  and  installation  sections to do so. In order to find where and which kernel is being used for your target, you can first find the recipes for the kernels used by different targets. You can also find this information
Read more

Customizable Embedded 3D Surround View Turn-Key Solution on Apalis iMX6 SoM

Image
Friday, June 23, 2017 Multi-Camera 3D Surround View systems are already established in many passenger vehicles. Demand for such systems now also grows for special vehicles like construction machines, harvesters, buses or military vehicles. In comparison to personal cars, these vehicles are produced in a rather moderate volume and have a high need for customization of the optical system. In this blog post, we present a customizable 3D Surround View Turn-Key solution based on the Toradex  Apalis iMX6 SoM  based on the NXP ®  i.MX 6 SoC, that can be quickly adapted concerning the specific requirements of the targeted vehicle. Please watch this video from Embedded World 2017 in Nuremberg, Germany, where this system solution was presented: We also explain how to get familiar with the topic of 3D surround view, evaluating the solution via a step-by-step approach and finally how to come to a product meeting your specific needs. TES 3D Surround View Stitching Library –
Read more