【I.MX6】imx6q使用openCL

imx6上有一个vivante gpu，通过gpu来加速浮点计算效果要好很多，但是通过学习过程中发现，这方面的资料少之又少，通过几天的折腾也算是能跑起来了，所以做一个记录。

下载openCL库

在使用openCL之前我们需要先安装openCL库，这些在nxp官网上提供了头文件和.so文件，我们直接下载放到编译器内即可。openCL库路径为：http://repository.timesys.com/buildsources/g/gpu-viv-bin-mx6q/ 我查了很多资料说只支持openCL1.1，我使用的是gpu-viv-bin-mx6q-1.1.0/ ，可以编译运行，别的版本没有试过。

安装openCL库

将下载的usr目录下lib目录和include目录下的文件分别放到编译器目录下的lib和include目录下即可，opt目录下有测试程序可以测试gpu是否正常运行。

编译openCL程序

测试程序在文章最后。编译命令为

arm-fsl-linux-gnueabi-g++ test.cpp -o test -lGAL -lOpenCL

编译过程中发现需要加-lGAL选项，不然编译报错，至于这个是干嘛的，为什么，没弄明白，如果你知道可以告诉我！

运行openCL程序

将可执行程序和.cl程序放到开发板直接运行即可。

---

测试程序：（忘记程序抄了哪的了，如有冒犯请告知）

（1）主程序test.cpp

#include <iostream>
#include <fstream>
#include <sstream>
#include <CL/cl.h>
 
const int ARRAY_SIZE = 1000;
 
//一、 选择OpenCL平台并创建一个上下文
cl_context CreateContext()
{
	cl_int errNum;
	cl_uint numPlatforms;
	cl_platform_id firstPlatformId;
	cl_context context = NULL;
 
	//选择可用的平台中的第一个
	errNum = clGetPlatformIDs(1, &firstPlatformId, &numPlatforms);
	if (errNum != CL_SUCCESS || numPlatforms <= 0)
	{
		std::cerr << "Failed to find any OpenCL platforms." << std::endl;
		return NULL;
	}
 
	//创建一个OpenCL上下文环境
	cl_context_properties contextProperties[] =
	{
		CL_CONTEXT_PLATFORM,
		(cl_context_properties)firstPlatformId,
		0
	};
	context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU,
		NULL, NULL, &errNum);
 
	return context;
}
 
 
//二、 创建设备并创建命令队列
cl_command_queue CreateCommandQueue(cl_context context, cl_device_id *device)
{
	cl_int errNum;
	cl_device_id *devices;
	cl_command_queue commandQueue = NULL;
	size_t deviceBufferSize = -1;
 
	// 获取设备缓冲区大小
	errNum = clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, NULL, &deviceBufferSize);
 
	if (deviceBufferSize <= 0)
	{
		std::cerr << "No devices available.";
		return NULL;
	}
 
	// 为设备分配缓存空间
	devices = new cl_device_id[deviceBufferSize / sizeof(cl_device_id)];
	errNum = clGetContextInfo(context, CL_CONTEXT_DEVICES, deviceBufferSize, devices, NULL);
 
	//选取可用设备中的第一个
	commandQueue = clCreateCommandQueue(context, devices[0], 0, NULL);
 
	*device = devices[0];
	delete[] devices;
	return commandQueue;
}
 
 
// 三、创建和构建程序对象
cl_program CreateProgram(cl_context context, cl_device_id device, const char* fileName)
{
	cl_int errNum;
	cl_program program;
 
	std::ifstream kernelFile(fileName, std::ios::in);
	if (!kernelFile.is_open())
	{
		std::cerr << "Failed to open file for reading: " << fileName << std::endl;
		return NULL;
	}
 
	std::ostringstream oss;
	oss << kernelFile.rdbuf();
 
	std::string srcStdStr = oss.str();
	const char *srcStr = srcStdStr.c_str();
	program = clCreateProgramWithSource(context, 1,
		(const char**)&srcStr,
		NULL, NULL);
 
	errNum = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
 
	return program;
}
 
//创建和构建程序对象
bool CreateMemObjects(cl_context context, cl_mem memObjects[3],
	float *a, float *b)
{
	memObjects[0] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
		sizeof(float) * ARRAY_SIZE, a, NULL);
	memObjects[1] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
		sizeof(float) * ARRAY_SIZE, b, NULL);
	memObjects[2] = clCreateBuffer(context, CL_MEM_READ_WRITE,
		sizeof(float) * ARRAY_SIZE, NULL, NULL);
	return true;
}
 
 
// 释放OpenCL资源
void Cleanup(cl_context context, cl_command_queue commandQueue,
	cl_program program, cl_kernel kernel, cl_mem memObjects[3])
{
	for (int i = 0; i < 3; i++)
	{
		if (memObjects[i] != 0)
			clReleaseMemObject(memObjects[i]);
	}
	if (commandQueue != 0)
		clReleaseCommandQueue(commandQueue);
 
	if (kernel != 0)
		clReleaseKernel(kernel);
 
	if (program != 0)
		clReleaseProgram(program);
 
	if (context != 0)
		clReleaseContext(context);
}
 
int main(int argc, char** argv)
{
	cl_context context = 0;
	cl_command_queue commandQueue = 0;
	cl_program program = 0;
	cl_device_id device = 0;
	cl_kernel kernel = 0;
	cl_mem memObjects[3] = { 0, 0, 0 };
	cl_int errNum;
 
	// 一、选择OpenCL平台并创建一个上下文
	context = CreateContext();
 
	// 二、 创建设备并创建命令队列
	commandQueue = CreateCommandQueue(context, &device);
 
	//创建和构建程序对象
	program = CreateProgram(context, device, "HelloWorld.cl");
 
	// 四、 创建OpenCL内核并分配内存空间
	kernel = clCreateKernel(program, "hello_kernel", NULL);
 
	//创建要处理的数据
	float result[ARRAY_SIZE];
	float a[ARRAY_SIZE];
	float b[ARRAY_SIZE];
	for (int i = 0; i < ARRAY_SIZE; i++)
	{
		a[i] = (float)i;
		b[i] = 1.0;
	}
 
	//创建内存对象
	if (!CreateMemObjects(context, memObjects, a, b))
	{
		Cleanup(context, commandQueue, program, kernel, memObjects);
		return 1;
	}
 
	// 五、 设置内核数据并执行内核
	errNum = clSetKernelArg(kernel, 0, sizeof(cl_mem), &memObjects[0]);
	errNum |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &memObjects[1]);
	errNum |= clSetKernelArg(kernel, 2, sizeof(cl_mem), &memObjects[2]);
 
	size_t globalWorkSize[1] = { ARRAY_SIZE };
	size_t localWorkSize[1] = { 1 };
 
	errNum = clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL,
		globalWorkSize, localWorkSize,
		0, NULL, NULL);
 
	// 六、 读取执行结果并释放OpenCL资源
	errNum = clEnqueueReadBuffer(commandQueue, memObjects[2], CL_TRUE,
		0, ARRAY_SIZE * sizeof(float), result,
		0, NULL, NULL);
 
	for (int i = 0; i < ARRAY_SIZE; i++)
	{
		std::cout << result[i] << " ";
	}
	std::cout << std::endl;
	std::cout << "Executed program succesfully." << std::endl;
	getchar();
	Cleanup(context, commandQueue, program, kernel, memObjects);
 
	return 0;
}

（2）核函数HelloWorld.cl

__kernel void hello_kernel(__global const float *a,
    __global const float *b,
    __global float *result)
{
    int gid = get_global_id(0);
        
    result[gid] = a[gid] + b[gid];
}