linux | 苹果OpenCL(提高应用软件如游戏、娱乐以及科研和医疗软件的运行速度和响应)

点击上方"蓝字"关注我们

01、引言

>>>
OpenCL 1.0 于 2008 年 11 月发布。

OpenCL 是为个人电脑、服务器、移动设备以及嵌入式设备的多核系统提供并行编程开发的底层 API。OpenCL 的编程语言类似于 C 语言。其可以用于包含 CPU、GPU 以及来自主流制造商如 NXP®、NVIDIA®、Intel®、AMD、IBM 等的处理器的异构平台。OpenCL 旨在提高应用软件如游戏、娱乐以及科研和医疗软件的运行速度和响应。

我们使用 Apalis iMX6Q 系统模块测试 OpenCL，对比两个应用 - 一个运行在 GPU 上，另一个则在 CPU。最后我们将分享本次测试的结果。

02、Toradex Embedded Linux 镜像中添加 OpenCL

>>>
假设你已经具有能够编译 Apalis iMX6 镜像的 OpenEmbedded 编译环境。你可以参考我们 OpenEmbedded (core) 【http://developer.toradex.com/knowledge-base/board-support-package/openembedded-(core)】文章。

为编译支持 OpenCL 以及相关库文件的嵌入式 Linux 镜像，需要采取以下步骤：

首先，修改下面目录中的文件。
/meta-toradex/recipes-fsl/packagegroups/packagegroup-fsl-tools-gpu.bbappend
添加如下内容：
SOC_TOOLS_GPU_append_mx6 = " \
    libopencl-mx6 \
    libgles-mx6 \
"
并在 local.conf 文件中添加 imx-gpu-viv
IMAGE_INSTALL_append = "imx-gpu-viv"
现在就可以编译镜像：
bitbake angstrom-lxde-image

03、GPU 和 CPU 代码

>>>

所有的代码可以从 GitHub【https://github.com/giobauermeister/OpenCL-test-apps】上下载。

我们使用数列求和应用作为基本的演示例程。第一部分代码运行在 GPU 上，第二部分则在 CPU 上。应用执行完毕后打印其所消耗的时间。使用 OpenCL 所需的头文件是 cl.h，位于文件系统的 /usr/include/CL 目录。链接程序所需的库文件是 libGAL.so 和 libOpenCL.so，位于 /usr/lib 目录。

为了计算消耗的时间，我们创建带分析功能的队列，在结束的时候获取分析的结果。

下面是 OpenCL 代码：

//************************************************************
// 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;
}

CPU 代码是简单的 C 程序，和上面一样计算同样的队列求和。为了计算消耗的时间，我们使用 time.h中的库。代码如下：

#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;
}

04、交叉编译应用

>>>
同一个 Makefile 可以用于交叉编译 GPU 和 CPU 应用。你需要注意下面的三个变量。根据你的系统做相应的调整：

ROOTFS_DIR -> Apalis iMX6 文件系统路径

APPNAME -> 应用的名字

TOOLCHAIN -> 交叉编译工具的路径
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)
在应用所在的目录中保持 Makefile 文件，然后运行 make。
将编译生成的文件复制到 Apalis iMX6 开发板上。

测试结果

在执行两个应用程序后，我们得到以下结果：
### 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
根据以上结果，我们可以很清楚地看到在 Apalis iMX6Q GPU 上使用 OpenCL 能够加速队列求和运算。

总结

>>>
借助 OpenCL，可以在不同设备从图形显卡到超级计算机以及嵌入式设备，运行代码。用户还可以进一步结合，例如在 OpenCV 中使用 OpenCL 提高计算机视觉的性能。

对于绝大多数嵌入式应用，Linux 是正确的选择。Linux 编译系统，例如 Buildroot 和 OpenEmbedded，能够创建定制化的 BSP，裁剪到任意的大小，并且提供丰富的应用和 SDK，从 gstreamer、Python 到 node.js 等。基于 OpenEmbedded/Yocto 的 Linux 是 Toradex 支持的默认发行版本，开发社区还提供多种开发语言环境和框架。

现在的 GUI 可以使用 Qt、HTML5 来开发，以至于有点难于选择。

你也可以使用 NDK 基于 C/C++ 和 C# 开发你的应用，使用 Qt 作为显示框架或者利用 Cordova 或者 React Native 框架用 Javascript 开发移动应用