OpenCL编程：主机端基本流程（平台、设备、上下文、程序、编译、内核、命令队列）

编写OpenCL程序的语法规则比较烦杂，要想构建一个复杂的OpenCL应用程序，就必须对诸如平台、设备、上下文、程序、内核以及命令队列等概念都有深入的理解和知识。

本文关注的主要是主机端编程，梳理一下主机端创建OpenCL程序的流程。为简便起见，本文未加入容错处理。

一、包含头文件

在使用OpenCL的API进行编程之前，首先要包含头文件，形式如下：

#ifdef __APPLE__
#include <OpenCL/opencl.h>
#else
#include <CL/cl.h>
#endif

二、获取平台信息

2.1 创建平台结构

编写平台程序一般分为两步。首先为一个或多个cl_platform_id结构分配内存空间。其次是调用函数clGetPlatformID来初始化这些数据结构。

函数原型如下：

cl_int clGetPlatformIDs(cl_uint num_entries, cl_platform_id *platforms, cl_uint *num_platforms);

实现代码：

cl_platform_id *platforms;
cl_uint num_platforms;
clGetPlatformIDs(5, NULL, &num_platforms);
platforms = (cl_platform_id*)malloc(sizeof(cl_platform_id) * num_platforms);
clGetPlatformIDs(num_platforms, platforms, NULL);

2.2 获取平台信息

创建了平台之后，我们本身并未得到任何关于平台的信息，这时需要调用一个名为clGetPlatformInfo的函数。

函数原型如下：

cl_int clGetPlatformInfo(cl_platform_id platform, cl_platform_info param_name, size_t param_value_size, void *param_value,

size_t *param_value_size_ret);

获取平台厂商的信息

实现代码：

char pform_vendor[40];
clGetPlatformInfo(platform[0], CL_PLATFORM_VENDOR, sizeof(pform_vendor), &pform_vendor, NULL);

三、访问安装设备

3.1 创建设备结构

我们需要创建一个cl_device_id结构来表示某个设备，调用clGetDeviceIDs。

函数原型如下：

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

实现代码：

 

cl_device_id *devices;
clGetDeviceIDs(platforms, CL_DEVICE_TYPE_ALL, 1, NULL, &num_devices);

3.2 获取设备信息

和创建平台一样，在创建了设备之后，我们也需要获取设备相关的信息。

函数原型如下：

cl_int clGetDeviceInfo(cl_device_id device, cl_device_info param_name, size_t param_value_size, void *param_value,

size_t *param_value_size_ret);

获取设备名

实现代码：

char name_data[40];
clGetDeviceInfo(device[0], CL_DEVICE_NAME, sizeof(name_data), name_data, NULL);

四、 通过上下文管理设备

上下文是命令队列创建的基础，而命令队列是主机和设备间通信的纽带。

4.1 创建上下文

OpenCL的上下文由cl_context结构来表示。

函数原型如下:

cl_context clCreateContext(const cl_context_properties *properties, cl_uint num_devices, const cl_device_id *devices,

(void CL_CALLBACK *notify_func)(...), void *user_data, cl_int *error);

cl_context clCreateContextFromType(const cl_context_properties *properties, cl_device_type device_type,

(void CL_CALLBACK *notify_func)(...), void *user_data, cl_int *error);

实现代码:

cl_context context;
cl_int err;
context=clCreateContext(NULL, 1, &device[0], NULL, NULL, &err);

4.2 获取上下文信息

同样，创建了上下文之后要获取信息，这和之前的情况非常类似。

函数原型如下：

clGetContextInfo(cl_context context, cl_context_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret);

获取上下文的引用计数

实现代码：

cl_uint ref_count;
clGetContextInfo(context, CL_CONTEXT_REFERENCE_COUNT, sizeof(ref_count), &ref_count, NULL);

五、将设备代码保存在程序中

5.1 创建程序

在OpenCL中， 一个程序由数据结构cl_program来表示。

函数原型如下：

clCreateProgramWithSource(cl_context context, cl_uint src_num, const char **src_strings, const size_t *src_sizes, cl_int *err_code);

clCreateProgramWithBinary(cl_context context, cl_uint num_devices, const cl_device_id *devices, const size_t *bin_sizes,

const unsigned char **bins, cl_int *bin_status, cl_int *err_code);

实现代码：

cl_program program;
FILE *program_handle;
char *program_buffer;
size_t program_size;
program_handle = fopen("kernel.cl", "r");
fseek(program_handle, 0 , SEEK_END);
program_size = ftell(program_handle);
rewind(program_handle);
program_buffer = (char*)malloc(program_size+1);
program_buffer[program_size] = '\0';
fread(program_buffer, sizeof(char), program_size, program_handle);
fclose(program_handle);
program = clCreatePrpgramWithSource(context, 1, (const char**)program_buffer, (const size_t*)&program_size, &err);

5.2 编译程序

很多厂商的编程框架并没有包含独立的OpenCL编译器。OpenCL标准也并没有对OpenCL编译器施加很大的压力，但是有一条必须遵守：每个编译器都必须能通过函数clBuildProgram进行访问。所以，包含OpenCL内核函数的.cl文件是运行时编译，也就是说编译器的调用运行是整个应用程序运行的一部分。

函数原型如下：

clBuildProgram(cl_program program, cl_uint num_devices, const cl_device_id *devices, const char* options,

(void CL_CALLBACK *notify_func)(...), void *user_data);

实现代码：

const char options[] = "-cl-finite-math-only -cl-no-signed-zeros";
clBuildProgram(program, 1, &device[0], options, NULL, NULL);

5.3 获取程序信息

一旦创建和编译好程序，可以通过调用clGetProgramInfo获得程序相关的数据结构的信息，例如上下文和目标设备。调用clGetProgramBuildInfo获取程序的编译信息。

函数原型如下：

clGetProgramInfo(cl_program program, cl_program_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret);

clGetProgramBuildInfo(cl_program program, cl_device_id device, cl_program_build_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret);

实现代码：

char* program_log;
size_t log_size;
clGetProgramBuildInfo(program, device[0], CL_PROGRAM_BUILD_LOG, 0 ,NULL, &log_size);
program_log = (char*)malloc((log_size+1)*sizeof(char));
clGetProgramBuildInfo(program, device[0], CL_PROGRAM_BUILD_LOG, log_size+1 ,program_log, NULL);

六、将函数打包为内核

在程序编译和链接之后，就需要将函数打包为名叫内核的数据结构。

6.1 创建内核

OpenCL标准定义了两个函数通过cl_program结构来创建cl_kernel结构。

函数原型如下：

clCreateKernelsInProgram(cl_program program, cl_uint num_kernels, cl_kernel *kernels, cl_uint *num_kernels_ret);

clCreateKernel(cl_program program, const char *kernel_name, cl_int *error);

实现代码：

 

char kernel_name[] = "vecAdd";
cl_kernel kernel;
kernel = clCreateKernel(program, kernel_name, &err);

notes:必须有一个名为vecAdd的内核函数才能创建成功。

 

6.2 获取内核信息

 

如果创建完内核之后，你需要获取内核相关信息，可以调用clGetKernelInfo。

函数原型如下：

clGetKernelInfo(cl_kernel kernel, cl_kernel_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret);

实现代码：

 

char kernel_name[20];
clGetKernelInfo(kernel, CL_KERNEL_FUNCTION_NAME, sizeof(kernel_name), kernel_name, NULL);

七、用命令队列保存内核

7.1 创建命令队列

在OpenCL中，命令队列用cl_command_queue结构来表示。

函数原型如下：

clCreateCommandQueue(cl_context context, cl_device_id device, cl_command_queue_properties properties, cl_int *err);

实现代码:

 

cl_command_queue queue;
queue = clCreateCommandQueue(context, device[0], CL_QUEUE_PROFILING_ENABLE , &err);

7.2 入列内核执行命令

clEnqueueNDRangeKernel通过命令队列向设备发送内核执行命令。

函数原型如下：

clEnqueueNDRangeKernel(cl_command_queue queue, cl_kernel kernel, cl_uint work_dims, const size_t *global_work_offset,

const size_t *global_work_size, const size_t *local_work_size, cl_uint num_events,

const cl_event *wait_list, cl_event *event);

实现代码：

 

cl_event kernel_event;
err = clEnqueueNDRangeKernel(queue, kernel, 1, globalWorkSize, localWorkSize, 0, NULL, &kernel_event);

本文跟随OpenCL程序的编写流程，介绍了所要用到的主要数据结构。后面会陆续介绍到数据传输和划分等主题。