OpenCL原理及使用

1、概述

OpenCL（Open Computing Language）是一个为异构平台编写程序的框架，此异构平台可由CPU、GPU、DSP、FPGA或其他类型的处理器与硬件加速器所组成。OpenCL由一门用于编写kernels（在OpenCL设备上运行的函数）的语言（基于C99）和一组用于定义并控制平台的API组成。OpenCL提供了基于任务分割和数据分割的并行计算机制。

OpenCL，它由两部分组成：

• 在主机处理器（Host）执行的运行时 API（主要运行在CPU中）
• 基于 C99 标准扩展的 OpenCL C 语言，用于编写在设备处理器（OpenCL device）运行的内核（kernel）代码（运行在GPU或CPU若其它device中）

注意，OpenCL的Host并不一定就是CPU，kernel也并不一定就是运行在GPU中，比如kernel也可以运行在CPU中，OpenCL只是一套框架、接口，需要厂商支持并实现对应接口，如果某厂商的CPU或者其它芯片也支持，都是可以运行的，但常规上，Host即CPU，kernel则运行在GPU中

OpenCL执行流程大致如下：

后面我们将围绕运行时API接口以及kernel两部分展开

2、OpenCL API

OpenCL的接口有很多，但它们的套路基本一样，本文列举几个关键接口，更多接口请参见：OpenCL 3.0 Reference Pages

2.1 获取平台

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

参数说明：

• num_entries，要获取的平台数量，如果 platforms 非空，则 num_entries 不能为 0

• platforms ， 返回获取的平台对象数组

• num_platforms， 用于查询返回可用的平台数目，num_platforms 可设为 NULL 忽略

• 返回值，正常执行返回 CL_SUCCESS

2.2 获取设备

获取平台下的 OpenCL 设备 Device，查询设备的硬件参数

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

参数说明：

• platform，clGetPlatformIDs 获取的 Platform ID
• device_type，获取 OpenCL Device 的类型
• num_entries，要获取的设备数量
• devices，返回获取的设备对象数组
• num_devices，返回平台连接 device_type 类型设备数目，可设为 NULL 忽略
• 返回值，正常执行返回 CL_SUCCESS

设备类型有多种，比如最基础的就有CPU和GPU

2.3 创建上下文

选择获取的设备，创建上下文 Context。使用的API是 clCreateContext。OpenCL 使用 Context 管理命令队列、程序内核、内存等资源对象。

cl_context clCreateContext(
    const cl_context_properties* properties,
    cl_uint num_devices,
    const cl_device_id* devices,
    void (CL_CALLBACK* pfn_notify)(const char* errinfo, const void* private_info, size_t cb, void* user_data),
    void* user_data,
    cl_int* errcode_ret);

参数说明：

• properties，上下文属性数组，每一项属性包含一个枚举常量，一个属性值，数组以 0 结尾。properties 指定了创建 Context 基于的 Platform，也可以设为 NULL，程序实现时自行选择 Platform。
• num_devices，devices 中指定的设备数。
• devices，clGetDeviceIDs 返回的设备对象数组。
• pfn_notify，注册回调函数，当 OpenCL 创建上下文失败时会执行回调函数。没有回调可设为 NULL
• user_data，传递给回调函数 pfn_notify 的指针参数，可设为 NULL
• errcode_ret，返回错误码，如果 errcode_ret 设为 NULL 不再返回错误码。
• 返回值，OpenCL 上下文成功创建时，返回创建的 cl_contex t对象，errcode_ret 返回CL_SUCCESS。创建失败时返回 NULL，errorcode_ret 返回错误码

2.4 创建program

通过 OpenCL C 源码字符串或二进制两种方式之一创建内核程序 Program（即读取kernel代码或对应的二进制数据，生成program）。编译 Program 生成二进制，检查编译错误并获取二进制代码。使用二进制代码创建 Program 能显著减少编译时间。

首先，我们看看如何用字符串来创建Program

cl_program clCreateProgramWithSource(cl_context context,
                                     cl_uint count,
                                     const char **strings,
                                     const size_t *lengths,
                                     cl_int *errcode_ret)

参数说明：

• context，有效的Context对象
• count，表示 strings 中字符串的个数
• strings，字符串数组指针，所有的字符串构成设备源代码
• lengths，表示 strings 每个字符串的长度。lengths 可以设为 NULL，字符串以 0 结尾自动计算长度。
• errcode_ret，返回错误码。errcode_ret 设为 NULL 则不再返回错误码
• 返回值：Program 对象成功创建时，返回创建的 cl_program 对象，errcode_ret 返回CL_SUCCESS。创建失败时返回 NULL，errorcode_ret 返回错误码

注意：strings是一个字符串指针数组，即意味着可以一次性传入多个kernel字符串，count即是strings这个数组的长度

用二进制数据来创建Program

cl_program clCreateProgramWithBinary(cl_context context,
                                     cl_uint num_devices,
                                     const cl_device_id *device_list,
                                     const size_t *lengths,
                                     const unsigned char **binaries,
                                     cl_int *binary_status,
                                     cl_int *errcode_ret)

参数说明：

• context，有效的 Context 对象。
• num_devices，device_list 中设备个数
• device_list，context 关联的设备数组。二进制需要载入 device_list 所列出的设备中，因此不能为NULL。
• lengths，binaries 数组中每个二进制文件的长度。
• binaries ， 二进制文件数组。对于设备device_list[i]，其程序二进制文件是 binaries[i]，文件的长度是 lengths[i]，三者一一对应。
• binary_status，返回每个设备对应的二进制是否成功加载。成功加载返回 CL_SUCCESS。binary_status可设为NULL以忽略。
• errcode_ret，返回错误码。errcode_ret 设为NULL则不再返回错误码。
• 返回值，Program 对象成功创建时，返回创建的 cl_program 对象，errcode_ret 返回CL_SUCCESS。创建失败时返回NULL，errorcode_ret 返回错误码。

接下来创建Program，编译生成可执行程序是clBuildProgram函数

cl_int clBuildProgram(cl_program program,
                      cl_uint num_devices,
                      const cl_device_id *device_list,
                      const char *options,
                      void (CL_CALLBACK *pfn_notify)( cl_program program, void *user_data),
                      void *user_data)

参数说明：

• program，创建的程序对象
• device_list，program 关联的设备对象数组。device_list 设为 NULL 时，为 program 关联的所有设备编译可执行程序。device_list 非空则仅为 device_list 中给出的设备编译可执行程序。

注意，编译kernel生成可执行程序是发生在clBuildProgram这个函数中。如果使用二进制数据生成program，那么节省的时间即是此函数执行的时间。另外，编译kernel生成的可执行程序是与device有关的，不同的device，可执行程序可能会不相同。

下面的代码展示了如何生成二进制program数据：

    cl_uint num_devices;
    cl_int err   = 0;
    // 1. 获取程序关联设备数
    err |= clGetProgramInfo(program, CL_PROGRAM_NUM_DEVICES, sizeof(cl_uint), &num_devices, NULL);
    LOGI("num_devices: %d", num_devices);
    // 2. 获取程序关联设备ID
    cl_device_id *p_devices = (cl_device_id *)malloc(sizeof(cl_device_id) * num_devices);
    err |= clGetProgramInfo(program, CL_PROGRAM_DEVICES, sizeof(cl_device_id) * num_devices,
                            p_devices, NULL);
    // 3. 获取设备程序二进制代码长度
    size_t *p_program_binary_sizes = (size_t *)malloc(sizeof(size_t) * num_devices);
    // 4. 获取设备程序二进制代码
    err |= clGetProgramInfo(program, CL_PROGRAM_BINARY_SIZES, sizeof(size_t) * num_devices,
                            p_program_binary_sizes, NULL);

    cl_uchar **p_program_binaries = (cl_uchar **)malloc(sizeof(cl_uchar *) * num_devices);
    for (cl_uint i = 0; i < num_devices; i++)
    {
        p_program_binaries[i] = (cl_uchar *)malloc(p_program_binary_sizes[i]);
        LOGI("Binary size for device %d=%zu\n", i, p_program_binary_sizes[i]);
    }
    err |= clGetProgramInfo(program, CL_PROGRAM_BINARIES, sizeof(cl_uchar *) * num_devices,
                            p_program_binaries, NULL);
    //保存二进制数据到文件中，后面就可以直接读文件获取二进制数据来生成program了

2.5 创建队列

为单个设备创建命令队列，使用的 API 是 clCreateCommandQueueWithProperties。操作命令入队后依据队列属性顺序或者乱序执行。

cl_command_queue 
clCreateCommandQueueWithProperties(cl_context context,
                                   cl_device_id device,
                                   const cl_queue_properties *properties,
                                   cl_int *errcode_ret)

参数说明：

• properties，命令队列属性数组，每一项属性包含一个枚举常量，一个属性值，数组以 0 结尾

注意，队列也是会与device关联的。到目前为止，context、program以及队列都会与device相关联，后面要执行kernel或读kernel的输出数据，都是通过队列。

2.6 创建内核对象

使用生成好的 Program 对象创建内核对象 kernel，类型为 cl_kernel

cl_kernel clCreateKernel(cl_program program,
                         const char *kernel_name,
                         cl_int *errcode_ret)

参数说明：

• program，已经生成可执行二进制的内核程序对象。必须是执行完clBuildProgram方法之后的program

2.7 数据

OpenCL传输给kernel的数据可以分成三类

• 标量数据，即普通值，如整形、浮点型之类的
• cl_mem，可理解为普通的buffer数据
• image，类型同样是cl_mem

buffer数据与image都需要使用特定接口创建、读写

首先来说明下如何创建buffer

cl_mem clCreateBuffer(cl_context context,
                        cl_mem_flags flags,
                        size_t size,
                        void *host_ptr,
                        cl_int *errcode_ret)

参数说明：

• flags，以组合 bit 位枚举常量的方式，指定 buffer 的分配和使用信息
• host_ptr，应用程序在 Host 端已经申请的内存空间指针

简单来说，此接口把Host端上的一个指针转变成kernel可用的buffer

再说明下如何读取buffer：

cl_int clEnqueueReadBuffer (cl_command_queue command_queue,
                            cl_mem buffer,
                            cl_bool blocking_read,
                            size_t offset,
                            size_t size,
                            void *ptr,
                            cl_uint num_events_in_wait_list,
                            const cl_event *event_wait_list,
                            cl_event *event)
cl_int clEnqueueWriteBuffer(cl_command_queue command_queue,
                            cl_mem buffer,
                            cl_bool blocking_write,
                            size_t offset,
                            size_t size,
                            const void *ptr,
                            cl_uint num_events_in_wait_list,
                            const cl_event *event_wait_list,
                            cl_event *event)

接口说明：

• clEnqueueReadBuffer ，从 buffer 对象读取数据到 Host 端内存
• clEnqueueWriteBuffer，将 Host 内存数据写入 buffer 对象

Image 对象封装了图像大小、图像格式、坐标模式、插值模式等多种信息。Image 对象在 Device 端可使用采样器 Sampler 方便地读取图像

cl_mem clCreateImage(cl_context context,
                     cl_mem_flags flags,
                     const cl_image_format *image_format,
                     const cl_image_desc *image_desc,
                     void *host_ptr,
                     cl_int *errcode_ret)

接口说明：

• flags，指定缓冲区的分配和使用信息，枚举类型与 clCreateBuffer 一致。
• image_format，图像格式，包括图像通道顺序和数据类型。
• image_desc，图像描述，包括图像类型、图像宽高和 pitch 等参数。
• host_ptr，host端内存地址，可用于初始化图像数据或设为 NULL。

cl_image_format 结构体包含*image_channel_order* 和image_channel_data_type **两个成员，举例来说，image_channel_order=CL_RGBA ， image_channel_data_type=CL_UNSIGNED_INT8表示图像每个像素包括 RGBA 四个通道，每个通道的数据都是 8 位无符号整数。

2.8 执行内核

设置 kernel 的任务网格尺寸，并执行内核。

提交内核执行命令，提交后 API 立即返回 Host。设备会按照设定的 work-item 网格尺寸启动内核函数执行。

cl_int clEnqueueNDRangeKernel(cl_command_queue command_queue,
                              cl_kernel kernel,
                              cl_uint work_dim,
                              const size_t *global_work_offset,
                              const size_t *global_work_size,
                              const size_t *local_work_size,
                              cl_uint num_events_in_wait_list,
                              const cl_event *event_wait_list,
                              cl_event *event)

参数说明：

• work_dim，work-item 的组织维度，0 < work_dim <= 3，全局 work-item 和工作组内work-item维度相同。
• global_work_offset，数组，表示 0~ work_dim-1 维全局工作项 ID 的偏移量。可设为 NULL，每个维度偏移量为 0。
• global_work_size，全局工作项尺寸数组，全局工作项总数为 global_work_size[0]*...* global_work_size[work_dim-1] 。
• local_work_size，工作组尺寸数组，工作组内工作项个数为local_work_size[0]*...* local_work_size[work_dim-1] 。

注意，global_work_size中的每一个维度数据必须可以整除local_work_size中对应的每一维数据

3、OpenCL 示例代码

接下来我们将使用opencl完成一次加法

首先是host代码：

    // 获取平台和设备信息
    cl_platform_id platform_id = NULL;
    cl_device_id device_id = NULL;
    cl_uint ret_num_devices;
    cl_uint ret_num_platforms;
    cl_int ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms);
    ret = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, &ret_num_devices);

    // 创建OpenCL上下文
    cl_context context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret);

    // 创建命令队列
    cl_command_queue command_queue = clCreateCommandQueueWithProperties(context, device_id, 0, &ret);

    // 准备内核源代码
    const char *kernel_filename = "kernel/vector_add.cl";
    FILE *fp = fopen(kernel_filename, "r");
    if (!fp) {
        fprintf(stderr, "Failed to load kernel file: %s\n", kernel_filename);
        exit(1);
    }
    
    char *source_str = (char*)malloc(MAX_SOURCE_SIZE);
    size_t source_size = fread(source_str, 1, MAX_SOURCE_SIZE, fp);
    fclose(fp);
    printf("source: %s\n", source_str);
    
    // 创建程序对象
    cl_program program = clCreateProgramWithSource(context, 1, (const char **)&source_str, 
                                                  (const size_t *)&source_size, &ret);
    free(source_str);

    // 构建程序
    ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);
    if (ret != CL_SUCCESS) {
        // 显示构建错误
        size_t log_size;
        clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
        char *log = (char *)malloc(log_size);
        clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, log_size, log, NULL);
        fprintf(stderr, "Error in kernel:\n%s\n", log);
        free(log);
        exit(1);
    }

    // 创建内核对象
    cl_kernel kernel = clCreateKernel(program, "vector_add", &ret);

    // 准备输入数据
    const int LIST_SIZE = 1024;
    float *a = (float *)malloc(sizeof(float) * LIST_SIZE);
    float *b = (float *)malloc(sizeof(float) * LIST_SIZE);
    float *c = (float *)malloc(sizeof(float) * LIST_SIZE);

    for (int i = 0; i < LIST_SIZE; i++) {
        a[i] = (float)i;
        b[i] = (float)(LIST_SIZE - i);
    }

    // 创建内存缓冲区
    cl_mem a_mem_obj = clCreateBuffer(context, CL_MEM_READ_ONLY, 
                                     LIST_SIZE * sizeof(float), NULL, &ret);
    cl_mem b_mem_obj = clCreateBuffer(context, CL_MEM_READ_ONLY, 
                                     LIST_SIZE * sizeof(float), NULL, &ret);
    cl_mem c_mem_obj = clCreateBuffer(context, CL_MEM_WRITE_ONLY, 
                                     LIST_SIZE * sizeof(float), NULL, &ret);

    // 将数据写入内存缓冲区
    ret = clEnqueueWriteBuffer(command_queue, a_mem_obj, CL_TRUE, 0,
                              LIST_SIZE * sizeof(float), a, 0, NULL, NULL);
    ret = clEnqueueWriteBuffer(command_queue, b_mem_obj, CL_TRUE, 0,
                              LIST_SIZE * sizeof(float), b, 0, NULL, NULL);

    // 设置内核参数
    ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&a_mem_obj);
    ret = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&b_mem_obj);
    ret = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&c_mem_obj);

    // 执行内核
    size_t global_item_size = LIST_SIZE;
    size_t local_item_size = 64;
    ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, 
                                &global_item_size, &local_item_size, 0, NULL, NULL);

    // 读取结果
    ret = clEnqueueReadBuffer(command_queue, c_mem_obj, CL_TRUE, 0,
                             LIST_SIZE * sizeof(float), c, 0, NULL, NULL);

    // 验证结果
    int correct = 1;
    for (int i = 0; i < LIST_SIZE; i++) {
        if (c[i] != a[i] + b[i]) {
            correct = 0;
            break;
        }
    }

    if (correct)
        printf("Vector add test PASSED!\n");
    else
        printf("Vector add test FAILED!\n");
    
    // 清理
    ret = clFlush(command_queue);
    ret = clFinish(command_queue);
    ret = clReleaseKernel(kernel);
    ret = clReleaseProgram(program);
    ret = clReleaseMemObject(a_mem_obj);
    ret = clReleaseMemObject(b_mem_obj);
    ret = clReleaseMemObject(c_mem_obj);
    ret = clReleaseCommandQueue(command_queue);
    ret = clReleaseContext(context);
    free(a);
    free(b);
    free(c);

kernel代码：

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

看起来，opencl好像非常简单，但实质上不是这样的。opencl复杂的地方有如下几个：

• 数据的映射，在推理引擎开发中，会有多个维度的数据，如NHWC等，而且一般我们会使用image作为核函数的数据类型，因为像gpu有纹理缓存、取image数据也会相对更加高效，但image只有宽高，多维数据如何映射？
• opencl的本质，是将工作分组、数据分组，然后并行计算，工作分组即是 clEnqueueNDRangeKernel 函数的分组参数，分组总数必须超过或等于数据分组总数，否则数据无法区分

这些内容我们下篇再说