OpenCL与CNN篇二：OpenCL基础API介绍

1. 本篇介绍几个OpenCL基础API，涉及平台、设备等初始化必备函数
2. 其次介绍几个关于缓冲区操作以及工作空间划分的API
3. 建议阅读参考书籍，我的推荐是《OpenCL编程指南》和《OpenCL异构计算》，尤其是后者从实践出发，更是适合上手。
4. 建议收藏这个API的参考网址，里面也OpenCL API的详细介绍OpenCL 2.1 Reference Pages
5. 本章节函数使用举例见系列篇三，这里不再重复；仅介绍API相关参数，以及自己的见解。

一、OpenCL平台API

平台API定义了宿主机程序发现OpenCL设备所用的函数以及这些函数的功能，另外还定义了为OpenCL
应用创建上下文的函数。包括平台、设备、上下文等相关函数。
不过为了方便起见这里将构建程序和创建内核也放在了此章节，即1.4、1.5、1.6。
注意1.7及之后为一些篇外话，不作为和1.1~1.6的连续内容。

1.1、获取平台clGetPlatformIDs()

cl_int clGetPlatformIDs(        cl_uint          num_entries,    //限制返回的平台数       cl_platform_id  *platforms,      //平台列表保存位置指针       cl_uint         *num_platforms   //平台个数保存位置指针        );

将num_entries和platforms分别设置为0和NULL可以查询可用的平台个数，保存在num_platforms中。返回的平台数可以用num_entries来限制，获取平台列表时这个参数要大于0并小于等于可用平台数。
平台这个概念应该是不同厂商的实现，比如我一开始安装了AMD APP SDK，运行程序只有1个可用平台；后来有想法就又安装了Intel的SDK，这时运行OpenCL有两个可用平台。区别是一个平台包含GPU和CPU，一个平台只包含CPU；而且打印设备信息，他们包含的CPU名称相同。

1.2、获取设备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        );

这个命令会得到与platform关联的eOpenCL设备列表。如果参数devices为NULL，则用num_devices得到可用设备数。类似平台函数用num_entries可以限制返回的设备数。
参数device_type用来指定计算设备类型，其极常见参数选择见下表。

cl_device_type 描述 CL_DEVICE_TYPE_CPU 作为宿主机处理器的OpenCL设备 CL_DEVICE_TYPE_GPU 作为GPU的OpenCL设备 CL_DEVICE_TYPE_ACCELERATOR OpenCL加速器(例如,IBM Cell Broaband) CL_DEVICE_TYPE_DEFAULT 默认设备 CL_DEVICE_TYPE_ALL 与相应平台关联的所有OpenCL设备

1.3、创建上下文clCreateContext()

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     );

第一参数着重说明一下，需要用确定的平台来构建，方法如下：
　　cl_context_properties props[3] =
　　{ CL_CONTEXT_PLATFORM , (cl_context_properties)platformIds, 0 };
第四个参数那么复杂就直接忽略吧，因为调用时那个位置填NULL。使用中user_data也输入NULL。最后一个参数errcode_ret用来记录错误代码，正常时errcode_ret==CL_SUCCESS。

1.4、创建程序对象clCreateProgramWithSource()

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

context 创建程序对象的上下文
count 没看到过介绍
const strings 将这个参数的所有字符，构成了创建程序对象的完整源代码
lengths 这个参数可以设置为NULL，这种情况下，则认为字符串以null终止的
errcode_ret 如果非NULL，函数返回错误代码将由这个参数返回
一般执行时，先从.cl文件中读取代码，然后将保存代码的字符串传递给strings

1.5、构建clBuildProgram()

cl_int clBuildProgram(       cl_program          program,       // 一个合法的程序对象       cl_uint             num_devices,   // 要构建程序对象的设备数       const cl_device_id *device_list,   // 如果为空则为context所有关联设备创建       const char         *options,       void(CL_CALLBACK   *pfn_notify)(cl_program program,void *user_data),       void               *user_data       );

函数的第四、五、六个参数设置为NULL即可，前三个参数根据代码段注释传递即可。深入学习的话关于第四个参数options，可以参考《OpenCL编程指南》6.2.2或者查看文章开头给的参考网页。

1.6、创建内核clCreateKernel()

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

到这里就比较容易了，只介绍一下kernel_name：创建内核对象的内核函数名。这是程序源代码__kernel关键字后面的内核函数名。参数errcode_ret的作用和上面的API中意义相同。

1.7、篇外话之先创建上下文再获取设备

• a、创建上下文clCreateContextFromType()
• b、从上下文中查询设备信息clGetContextInfo()

cl_context clCreateContextFromType(     const cl_context_properties *properties,     cl_device_type               device_type,     void (CL_CALLBACK *pfn_notify)(const char *,const void *p,size_t, void *),     void                        *user_data,     cl_int                      *errcode_ret                );cl_int clCreatContextInfo(       cl_Context      context,       cl_Context_info param_name,       size_t          param_value_size,       void           *param_value,       size_t         *param_value_size_ret       )

和clCreateContext()类似，不过这里不需要事先指定设备列表，使用实例：
　　platform = platformIds[0];
　　cl_context_properties props[3] =
　　{ CL_CONTEXT_PLATFORM,(cl_context_properties)platform, 0 };
　　context = clCreateContextFromType( props, CL_DEVICE_TYPE_CPU, NULL, NULL, &errNum);
之后需要在上下文中查询设备clGetContextInfo()，获取到设备列表之后操作就和之前相同了，使用实例：
　　cl_uint numDevices ;
　　cl_device_id *deviceIds;
　　errNum = clGetContextInfo( context, CL_CONTEXT_DEVICES, 0, NULL, &numDevices);
　　deviceIds =(cl_device_id*)malloc(sizeof(cl_device_id)*numDevices);
　　errNum = clGetContextInfo( context, CL_CONTEXT_DEVICES, numDevices, deviceIds, NULL);
把这一个篇外话，是因为这有一个不可取之处，这样搞会有很奇怪的问题。首先在我的电脑上如果是用clCreateContextFromType创建上下文设备选择CL_DEVICE_TYPE_CPU；在调用clGetContextInfo获取的numDevices=4，这是我CPU的线程数，并不是CPU个数。这样本来我还能接受，因为只选择deviceIds[0]的话程序可以正常执行；但如果打印设备信息，deviceIds[0]的信息能正常获取并且printf出来，但是获取deviceIds[1]时就开始报错了，简直。。。。。
把这一个篇外话，只是为了在实践中理解上下文和平台设备之间的关系，并不是推荐使用。

二、OpenCL运行时API

这些API管理上下文来创建命令队列以及运行时发生的其他操作。包括创建读写缓冲区、设置执行内核等

2.1、设置内核参数clSetKernelArg()

cl_int clSetKernelArg(       cl_kernel   kernel,       cl_uint     arg_index,       size_t     *arg_size,       const void *arg_value       );

kernel 一个合法的内核对象；arg_index 内核参数的索引；arg_size 参数的大小；arg_value 传入内核函数的参数的一个指针。这里参数的意思看看就好，重要的是看例子学会使用。设置好kernel参数之后通过调用下文的clEnqueueNDRangeKernel来执行。
值得说明的是如果你要多次调用执行同一个内核，那么可以只设置一次内核。例如我写卷积神经网络，在程序开始设置内核参数，之后训练过程中只需把新的图像写入到图像缓冲区，直接调用clEnqueueNDRangeKernel执行内核就好。

2.2、执行内核clEnqueueNDRangeKernel()

cl_int clEnqueueNDRangeKernel(       cl_command_queue commad_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  *events_in_wait_list,       cl_event        *event       );

commend_queue 内核的执行需要提交到命令队列中
kernel 要执行的内核名称
work_dim 指定新建的work-item的纬度，这里我假设你已经知道了OpenCL的内核与工作项之间的关系。
global_work_offset 为work-item提供全局ID，该参数可以不从0开始，但一般设置为0或者NULL
global_work_size 指定NDRange中每维work-item的数量，不可为空
local_work_size 指定workgroup中每维work-item的数量，可以设置为NULL让系统自动设置
num_events_in_wait_list、events_in_wait_list、event 这是OpenCL高级一点的操作，用于记录此事件或者需要等待的其他事件，可以用来规划不同事件执行顺序；如果不用可以分别设置为0、NULL、NULL。
需要注意的是global_work_size指向的数组大小要和work_dim相等，global_work_offset和local_work_size不为NULL时也是同样要求，不一样会发生什么我也没有测试，不过正常的程序这里该是一样的。

三、缓冲区的创建、写入和读取

3.1、flush和finish命令

这两个命令在命令队列中值两种不同类型的barrier操作。clFinish()函数阻塞直到命令队列中所有命令完成。clFlush()阻塞指导命令队列中的所有命令被移出队列，这意味着这些命令已经准备就绪但无法保证执行完毕。
cl_int clFlush(cl_command_queue command_queue);
cl_int clFinish(cl_command_queue command_queue);

3.2、创建缓冲区clCreateBuffer()

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

context 一个合法的上下文，为这个上下文分配缓冲区
size 所分配缓冲区的大小（字节数）
host_ptr 这个指针在clCreateBuffer如何使用有flags参数确定。host_ptr指向的数据大小应大于等于size
flags 用于指定关于缓冲区创建的分配和使用信息。其部分取值见下表。

cl_mem_flags 描述 CL_MEM_READ_WRITE 指定内存对象将由内核读写；默认为此模式 CL_MEM_WRITE_ONLY 指定内存对象由内核写，但不能读。 CL_MEM_READ_ONLY 指定内存对象由内核读，但不能写。 CL_MEM_USE_HOST_PTR 只有当host_ptr为非NULL时，这个标志合法；使用host_ptr引用的内存作为内存对象的存储位 CL_MEM_ALLOC_HOST_PTR 指定缓冲区应当在宿主机可访问的内存中分配。不可与USE_HOST_PTR同时使用 CL_MEM_COPY_HOST_PTR 表示希望OpenCL实现分配内存对象的内存，并从host_ptr引用复制数据。

3.3、读写缓冲区

cl_int clEnqueueWriteBuffer(       cl_command_queue commad_queue,       cl_mem           buffer,       cl_bool          blocking_write,       size_t           offset,       size_t           cb,       void            *ptr,       cl_uint          num_events_in_wait_list,       const cl_event  *events_in_wait_list,       cl_event        *event       );cl_int clEnqueueReadBuffer(       cl_command_queue commad_queue,       cl_mem           buffer,       cl_bool          blocking_read,       size_t           offset,       size_t           cb,       void            *ptr,       cl_uint          num_events_in_wait_list,       const cl_event  *events_in_wait_list,       cl_event        *event          );

command_queue 这是一个命令队列,读写命令将在这个队列中排队
buffer 一个合法的缓冲区对象（数据对这里读写）
blocking_read 如果设置为CL_TRUE，则命令阻塞,直至ptr读写数据完成
offset 缓冲区对象的读写数据的起始偏移量（字节数）
cb 对缓冲区读写的字节数
ptr 宿主机内存的一个指针，写入缓冲区的数据从哪里来 / 或者从缓冲区读数据写入哪里
关于读写的blocking_write的多一些说明：
　　如果blocking_write为CL_TRUE,
　　　　则OpenCL实现将复制ptr引用的数据，并在命令队列中对写操作进行排队。
　　　　在clEnqueueWriteBuffer调用返回后，由ptr指向的内存可以被应用程序重用。
　　如果blocking_write为CL_FALSE,
　　　　则OpenCL实现将使用ptr执行非阻塞写操作。 由于写是非阻塞的，实现可以立即返回。
　　　　ptr指向的内存在调用返回后不能被应用程序重用。
　　　　event参数返回一个事件对象，可以用来查询write命令的执行状态。
　　　　当写命令完成后，ptr指向的内存可以被应用程序重新使用

四、图像创建读写

图像类型的数据我还没有用过，不过好像挺有用的；在此只给出API的介绍。

4.1、图像格式

typedef struct _cl_image_format{    cl_channel_order image_channel_order;    cl_channel_type  iamge_channel_date_type; }   cl_image_format ;

The image format describes how the data will be stored in memory
使用示例:
cl_image_format format;
format.image_channel_order = CL_R; // single channel
format.image_channel_data_type = CL_FLOAT; // float data type

4.2、图像创建

cl_mem clCreateImage2D(       cl_context             context,       cl_mem_flags           flags,       const cl_image_format *image_format       size_t                 image_with,       size_t                 image_height,       size_t                 image_row_pitch,       void                  *host_ptr,       cl_int                *errcode_ret          );cl_mem clCreateImage3D(       cl_context             context,       cl_mem_flags           flags,       const cl_image_format *image_format       size_t                 image_with,       size_t                 image_height,       size_t                 image_depth,       size_t                 image_row_pitch,       size_t                *image_slice_pitch,       void                  *host_ptr,       cl_int                *errcode_ret          );

context 创建图像对象的上下文
flags 其合法枚举由cl_mem_flags定义
image_format 描述通道次序和图像通道数据类型
image_with,image_height,image_depth 图像的长宽深
image_row_pitch 如果host_ptr不为NULL,这个值指定图像中各行的字节数;为0采取默认值
image_slice_pitch 如果host_ptr不为NULL,这个值指定图像中各个切片的字节数;为0采取默认值
host_ptr 内存中线性布局的图像缓冲区指针
errcode_ret 如果为非NULL,函数返回错误码
使用示例:
cl_mem d_inputImage = clCreateImage2D(context, 0, &format, imageWidth, imageHeight, 0, NULL, &errNum);

4.3、图像读写

cl_int clEnqueueWriteImage(       cl_command_queue commad_queue,       cl_mem           image,       cl_bool          blocking_read,       const size_t     origin[3],       const size_t     region[3],       size_t           row_pitch       size_t           slice_pitch,       void            *ptr,       cl_uint          num_events_in_wait_list       const cl_event  *event_wait_list,       cl_event        *event          );

commad_queue 写入命令将放入这个队列
iamge 这是一个合法的图像对象
blocking_read 如果设置为CL_TRUE,则clEnqueueReadImage阻塞,直到数据读入ptr
origin 要写入相对图像原点的(x,y,z)整数坐标,对于二维z=0
region 要写入区域的(宽,高,深),对于二维z=1
row_pitch 图像中各行字节数,默认为image_with*(byte_per_pixel)
slice_pitch 三维图像中各切片的字节数image_height*row_pitch
ptr 这个指针指向源数据的宿主机内存
使用示例:
errNum = clEnqueueWriteImage(queue, d_inputImage, CL_FALSE, origin, region, 0, 0, inputImage, 0, NULL, NULL);

cl_int clEnqueueReadImage(       cl_command_queue commad_queue,       cl_mem           image,       cl_bool          blocking_read,       const size_t     origin[3],       const size_t     region[3],       size_t           row_pitch       size_t           slice_pitch,       void            *ptr,       cl_uint          num_events_in_wait_list       const cl_event  *event_wait_list,       cl_event        *event        );

commad_queue 读取命令将放入这个队列
iamge 这是一个合法的图像对象
blocking_read 如果设置为CL_TRUE,则clEnqueueReadImage阻塞,直到数据读入ptr
origin 要读取相对原图像原点的(x,y,z)整数坐标,对于二维z=0
region 要读取区域的(宽,高,深),对于二维z=1
row_pitch 图像中各行字节数,默认为image_with*(byte_per_pixel)
slice_pitch 三维图像中各切片的字节数image_height*row_pitch
ptr 这个指针指向写入数据的宿主机内存
使用示例:
errNum = clEnqueueReadImage(queue, d_outputImage, CL_TRUE, origin, region, 0, 0, outputImage, 0, NULL, NULL);

4.4、篇外话

1.要实现图像类数据的使用，好像还需要声明一个采样器（cl_sampler），具体怎么操作我还没有试验过。
2.在内核中读写图像好像需要使用固定的函数read_imagef()和write_imagef()；具体解释参考OpenCL Reference Pages =>OpenCL Compiler =>Built-in Functions => Image Functions =>read_imagef / write_imagef 。