背景

最近在学OpenCL，用的书是《OpenCL异构并行编程实战》。怎么说呢，感觉这本书比较迷，讲的很乱，跟著看完的话可能学不到什么。好在之前学过CUDA，勉强能够理解并行计算的思路。因此这里写了一下用显卡进行RGB2Gray的程序，也算是自己能够使用OpenCL进行一些简单的显卡计算吧。

平台：Ubuntu18.04 + CUDA10.0 + MX150 + OpenCL 1.2

目标：从磁碟读取一张图片，送入OpenCL设备进行RGB转灰度图，再拷贝回来显示

使用OpenCL进行计算的流程

无论是OpenCL，还是CUDA，当利用显卡计算时，都需要经历如下步骤：

1. 设备初始化 2. 准备主机端数据(分配主机端内存+获取数据) 3. 分配设备端内存 4. 将主机端数据拷贝到设备端 5. 设备启动内核函数，进行运算，将结果写到设备端内存 6. 将设备端结果拷贝回主机端 7. 读取主机端内存，进行后续处理 8. 释放资源

对于OpenCL，具体的步骤是：

1. 设备初始化(获取平台和设备id，创建上下文和命令队列) 2. 编写并编译内核 3. 准备主机端数据并传入设备(准备主机端数据，创建设备端缓冲对象，传入数据) 4. 启动内核函数(传递参数，启动内核) 5. 将结果拷贝回主机端 6. 后续处理 7. 释放资源

设备初始化

获取平台id

cl_int error;
cl_platform_id platform;

error = clGetPlatformIDs(1, &platform, NULL); //获取平台id

这里clGetPlatformIDs的函数原型为：

clGetPlatformIDs(cl_uint /* num_entries */,
cl_platform_id * /* platforms */,
cl_uint * /* num_platforms */)

本例是只读取一个平台。实际上可能有多个平台。因此实际上更为通用的使用方法为：

cl_uint numOfPlatforms;
cl_int error;

error = clGetPlatformIDs(0, NULL, &numOfPlatforms); //获取平台数量
if(error != CL_SUCCESS)
{
perror("Cannot get platform ids");
exit(1);
}
cl_platform_id *platforms = (cl_platform_id*)alloca(sizeof(cl_platform_id)*numOfPlatforms);
error = clGetPlatformIDs(numOfPlatforms, platforms, NULL); //获取cl_platform_id实体

这里我们看到clGetPlatformIDs()函数被调用了两次，第一次是获取平台的数量numOfPlatforms，第二次是获取cl_platform_id的实体。OpenCL里有很多类似的API。

本例是为了偷个懒，只使用第一个平台。

获取设备id

获取完平台id后，要获取设备id。这里一样偷懒，只使用平台上第一个设备。实际上一个平台上可能有多个OpenCL设备(CPU、GPU)。

cl_device_id device;

error = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL); //获取设备id

函数原型为：

clGetDeviceIDs(cl_platform_id /* platform */,
cl_device_type /* device_type */,
cl_uint /* num_entries */,
cl_device_id * /* devices */,
cl_uint * /* num_devices */)

创建设备上下文

OpenCL使用上下文来管理设备，因此不论进行什么操作，都需要先创建上下文：

context = clCreateContext(NULL, 1, &device, NULL, NULL, &error); //创建上下文

创建命令队列

命令队列和CUDA中的流的概念相对应。主机程序在设备上创建命令队列，并向命令队列中压入操作(数据传输、内核执行)。 一个设备上可以创建多个命令队列。

与CUDA略有区别的是，OpenCL除了顺序执行命令队列中的任务之外，还可以不按顺序执行(创建队列时传入CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE标志位)。

创建命令队列时，需要传入设备id和上下文

cQ = clCreateCommandQueue(context, device, NULL, &error);

函数原型：

clCreateCommandQueue(cl_context /* context */,
cl_device_id /* device */,
cl_command_queue_properties /* properties */,
cl_int * /* errcode_ret */)

编译内核

与CUDA不同的是，OpenCL需要在程序中显式编译内核。

内核函数

OpenCL工程一般将内核放在单独的*.cl文件中。本例的内核为：

/*rgb2gray.cl*/
__kernel void kernel_rgb2gray(__global unsigned char * rgbImage,
__global unsigned char * grayImage,
__global unsigned * const p_height,
__global unsigned * const p_width)
{
int x = get_global_id(0);
int y = get_global_id(1);
int height = *p_height;
int width = *p_width;
if(x < width && y < height)
{
int index = y * width + x;
grayImage[index] = 0.299f*rgbImage[index*3] +
0.587f*rgbImage[index*3+1] +
0.114f*rgbImage[index*3+2];
}
}

程序很简单，每个工作项(对应于CUDA中的线程)读取一个像素，计算灰度值，写回该像素。

由于图像尺寸不定，因此工作项不可能刚好和图像尺寸相对匹配，因此需要if(x < width && y < height)来分支。height和width则是由主机程序启动内核时传入的。

编译内核

OpenCL需要显示编译内核，其流程为：

读取内核文件->创建program对象->编译程序->创建内核

读取内核文件的代码如下：

const char *file_names[] = {"rgb2gray.cl"}; //待编译的内核文件
 const int NUMBER_OF_FILES = 1;
char* buffer[NUMBER_OF_FILES];
size_t sizes[NUMBER_OF_FILES];

loadProgramSource(file_names, NUMBER_OF_FILES, buffer, sizes); //读取内核文件文本
 program = clCreateProgramWithSource(context, NUMBER_OF_FILES, (const char**)buffer, sizes, &error); //创建program对象
 error = clBuildProgram(program, 1, &device, NULL, NULL, NULL); //编译程序

if(error != CL_SUCCESS) {
// If theres an error whilst building the program, dump the log
 size_t log_size;
clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
char *program_log = (char*) malloc(log_size+1);
program_log[log_size] =

相关文章