GPU加速——OpenCL学习与实践

原文首发于公众号「３D视觉工坊」：GPU——OpenCL学习与实践;

参考刘文志等所著《OpenCL异构并行计算》，结合自身实践所做的总结，在此，特别感谢蒋工给予的指导。由于作者认知水平有限，文中如有不到的地方，欢迎大家批评指正。

前言

由于CUDA完美地结合了C语言的指针抽象，NVIDIA不断升级其CUDA计算平台，CUDA获得了大量科学计算人员的认可，已经成为目前世界上使用最广泛的并行计算平台。通过CUDA，NVIDIA成功打破了Intel在超算市场上的绝对主导地位。在今天，大多数大中小型超算中心中都有GPU的身影。

由于CUDA由NIVIDA一家设计，并未被Intel和AMD等接受，因此目前使用CUDA编写的程序只支持NVIDA GPU，而OpenCL的出现解决了这一问题。

OpenCL全称为Open Computing Language(开放计算语言)，先由Apple设计，后来交由Khronos Group维护，是异构平台并行编程的开放标准，也是一个编程框架。Khronos Group是一个非盈利性技术组织，维护着多个开放的工业标准，并且得到了业界的广泛支持。OpenCL的设计借鉴了CUDA的成功经验，并尽可能地支持多核CPU、GPU或其他加速器。OpenCL不但支持数据并行，还支持任务并行。同时OpenCL内建了多GPU并行的支持。这使得OpenCL的应用范围比CUDA广。为了能适用于一些更低端的嵌入式设备（如DSP+单片机这种环境），OpenCL API基于纯C语言进行编写，所以OpenCL API的函数名比较长，参数也比较多（因为不支持函数重载），因此函数名相对难以熟记。不过，借助像Xcode、Visual Studio等现代化的集成开发环境，利用代码智能感知自动补全，其实开发人员也不需刻意去死背OpenCL的API。-- 引自《OpenCL 异构并行计算》

一 OpenCL的执行流程

1）查询平台

int getPlatform(cl_platform_id &platform)
{
    platform = NULL;//the chosen platform   cl_uint numPlatforms;//the NO. of platforms
 cl_int    status = clGetPlatformIDs(0, NULL, &numPlatforms);
    if (status != CL_SUCCESS)
    {
        cout<<"Error: Getting platforms!"<

2）查询设备

cl_device_id *getCl_device_id(cl_platform_id &platform)
{
    cl_uint numDevices = 0;
    cl_device_id *devices=NULL;
    // TODO
 cl_int    status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, NULL, &numDevices);
    //cl_int status;
 //numDevices = 1; cout << "Devices number:\t" << numDevices << endl;
    if (numDevices > 0) //GPU available.
 {
        devices = (cl_device_id*)malloc(numDevices * sizeof(cl_device_id));
        status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, numDevices, devices, NULL);
    }
    return devices;
}

3）创建上下文

mContext = clCreateContext(NULL,1, mDevices,NULL,NULL,&error);

4）创建命令队列

cl_command_queue commandQueue = clCreateCommandQueue(mContext, mDevices[0], 0, NULL);

5）创建内存对象

mpEcoMems[0] = clCreateBuffer(mContext, CL_MEM_READ_WRITE|CL_MEM_ALLOC_HOST_PTR, (FILTER_ROW_NUM) * sizeof(unsigned char),NULL, &error);
CHECK_CL_ERROR(msMemName[0]);

6）创建程序对象

此处有两种形式，一种是使用.cl文件（源码方式），便于调试。
static cl_program mProgramCl=0;  mProgramCl = clCreateProgramWithSource(mContext, 1, &source, sourceSize, &error);  另一种形式使用.cl文件编译成二进制后的二进制文件，速率较快，但不方便调试。
static cl_program mProgramCl=0;  mProgramCl = clCreateProgramWithBinary(mContext, 1, mDevices, &binary_pro_len,(const unsigned char**)&binary_pro, &binary_status, &error);

7）编译程序对象

编译程序即可。

8）创建内核对象

static cl_kernel mpEcoKernels[KERNEL_NUM];

9）设置内核参数

调用函数clSetKernelArg()即可。

10）执行内核

cl_event enentPoint;
size_t global_work_size[1] = {1};
status = clEnqueueNDRangeKernel(commandQueue, mpEcoKernels[0], 1, NULL, global_work_size, NULL, 0, NULL, &enentPoint);

二 创建命令队列

在OpenCL上下文中，有内存、程序和内核对象，对这些对象的操作就需要使用命令队列。一条命令就是主机发送给设备的一条消息，用来告诉设备执行一个操作。这个操作包含主机与设备间、设备内的数据拷贝与内核执行。命令提交到命令队列中，命令队列把需要执行的命令发送给设备。需要注意的是，每条命令队列只能关联一个设备，如果要同时使用多个设备，则需要创建多个命令队列，每个名列队列关联到一个设备，如下图所示。

三 工作项布局函数

理解两个概念：工作项与工作组。

某个学校高一的年级，这个年级当中会有多个班级，我们假设班级个数为8。高一年级都有计算机课程，会依次去计算机机房里上机，计算机机房里会有电脑，我们假设电脑数为32。每个在机房里的同学根据机房里黑板上老师布置的任务，都在完成属于自己的任务。

对于这样一个场景中的事物与OpenCL中几个概念的类比为：工作项就好比每位同学，工作组就好比一个班级，多个同学组成一个班级，多个工作项也组成一个工作组；机房里的电脑就好比处理单元，机房就好比计算单元。多个类似机房的计算单元构成了一个OpenCL设备。

我们以核心函数来体会OpenCL中的工作项与工作组的用法。

核心函数1： clEnqueueNDRangeKernel()

1）参数command_queue为提交内核执行任务的命令队列，命令队列创建时关联了指定的设备，command_queue关联的设备就是最后执行内核函数的设备。

2）参数)kernel为在设备上执行的内核函数。

3）参数work_dim制定设备上执行内核函数的全局工作项的维度。最为为值1，最大值为CL_DEVICE_MAX_WORK_TIME_DIMENSIONS。

4）参数global_work_offset为全局工作项ID的偏移量。如果global_work_offset为NULL，则偏移量为0。在目前的大多数设备上，此参数必须设置为NULL。

5）参数global_work_size指定全局工作项的大小。

6）参数local_work_size为一个工作组内工作项的大小。

关于参数work_dim、global_work_offset、global_work_size和local_work_size的用法，后面我会再进一步介绍。

7）参数num_events_in_wait_list和event_wait_list指定了在执行内核操作之前，需要等待num_events_in_wait_list各event_wait_list中的事件执行完成。

8）参数event指向这个命令生成的一个事件对象。后续的命令或主机可以使用这个事件的状态来控制其他操作。

四 缓冲区对象映射

可以使用如下函数来映射缓冲区对象到主机内存区域：

void *clEnqueueMapBuffer(
   cl_command_queue command_queue, //为一个有效的主机命令队列
  cl_mem buffer,
   cl_bool blocking_map,
//表明此映射是阻塞的还是非阻塞的。如果blocking_map为//CL_TRUE，即表示映射命令是阻塞的，直到映射完成，函数才会返回。应用可以用返回的指针访问所映射区域的内容；如果blocking_map为CL_FALSE，即映射为非阻塞的，直到映射命令完成后才能使用返回的指针。
 cl_map_flags map_flags, //用于描述映射区域的状态
 size_t offset,  //所要映射的区域在缓冲区对象中的偏移量，单位为字节
 size_t size,  //所要映射的区域在缓冲区对象中的大小，单位为字节
 cl_uint num_events_in_wait_list,
   const cl_event *event_wait_list,
    cl_event *event,  //参数event会返回一个事件对象，可用来查询映射命令的执行情况。
 //在映射完成后，应用才可以使用返回的指针访问映射区域的内容。
 cl_int *errcode_ret)

示例demo：将GPU上的数据映射到CPU内存，再将CPU上的内存映射回GPU。

//GPU——>CPU

static void* eco_map_mem(cl_command_queue &command_queue, int mem_id, cl_bool blocking_map, cl_map_flags map_flag)
{
   if(mem_id<0 || mem_id >= MEM_NUM)
   {
      LOCAL_COUT<

}

//CPU——>GPU

static int eco_unmap_mem(cl_command_queue &command_queue,int mem_id, void*ptr)
{
   if(mem_id<0 || mem_id >= MEM_NUM || !ptr)
   {
      LOCAL_COUT<

五 OpenCL事件机制

以一个简单示例来说明：

//我们这里用evt1来监测对src1MemObj做数据传输的命令执行状态 cl_event evt1,evt2;
ret =clEnqueueWriteBuffer(command_queue,src1MemObj，CL_FALSE,
0,contentLength,pHostBuffer,0,NULL,&evt1);
If(ret !=CL_SUCCESS)
{
puts(“Data1 transfer failed”);
Goto FINISH;
}

ret=clEnqueueWriterBuffer(command_queue,src2MemObj,CL_TRUE,
                          0,contentLength,pHostBuffer,02NULL,&evt1,&evt2);

if(ret !=CL_SUCCESS)
{
puts(“Data2 transfer failed”);
Goto FINISH;
}
clReleaseEvent(evt1);
clReleaseEvent(evt2);

如果我们的需求是在某个命令完成之前不想让当前主机端的线程继续往下执行该怎么办呢？如果使用clFinish函数，那么主机端的线程会被一直挂起，直到命令队列中所有命令全都执行完了之后才能返回操作。而如果我们仅仅只是等某个命令执行完成，就可以使用clWaitForEvents函数接口，其声明如下：

cl_int clWaitForEvents (cl_uint num_events, const cl_event *event_list)

这个函数会将主机端的线程挂起，直到event_list中的所有事件全部都完成。第一个参数num_events指定事件列表中一共有多少个事件需要等待完成。下面我们基于上述代码，在clReleaseEvent(evt1);上面添加如下代码：（绿色标注部分）

//我们这里用evt1来监测对src1MemObj做数据传输的命令执行状态 cl_event evt1,evt2;
ret =clEnqueueWriteBuffer(command_queue,src1MemObj，CL_FALSE,
0,contentLength,pHostBuffer,0,NULL,&evt1);
If(ret !=CL_SUCCESS)
{
puts(“Data1 transfer failed”);
Goto FINISH;
}

ret=clEnqueueWriterBuffer(command_queue,src2MemObj,CL_TRUE,
                          0,contentLength,pHostBuffer,02NULL,&evt1,&evt2);

if(ret !=CL_SUCCESS)
{
puts(“Data2 transfer failed”);
Goto FINISH;
}
struct timeval,tsBegin,tsEnd;
gettimeofday(&tsBegin,NULL);
clWaitForEvents(1,&evt2);
gettimeofday(&tsEnd,NULL);
long duration
        =1000000L*(tsEnd.tv_sec-tsBegin.tv_sec)+(tsEnd.tv_usec-tsBegin.tv_usec);

printf(“Wait time spent:%1dus\n”,duration);

clReleaseEvent(evt1);
clReleaseEvent(evt2);

六 原子操作

原子操作在传统并行计算上也是用得非常多的技巧。例如，对一些同步原语(synchronization primitive)的实现都可能会用到原子操作。最常见的就是多个线程如果要对同一存储地址的内容进行更新，就要用到原子操作进行访存。例如，我们要对一个大数组进行求和操作，倘若我们是在一个具有双核的处理器上执行，那么我们可能会将一个核的线程执行前一半求和，另一个核上的线程执行后一半，最后将这两个结果相加。如果我们的实现是把最终结果存放在一个全局变量里，这个变量的地址对于这两个线程而言都是可获得的。

那么一个线程做完一半求和之后用原子的加法操作对这个全局变量进行一次求和更新，这样，当另一个线程也用原子操作更新这个全局变量时结果是确定的。原子操作往往会对总线做

一次锁步操作，让当前总线上的访存操作能按次序进行。同时，又会刷新当前Cache，使得任一线程对全局变量使用了原子操作之后，其他所有线程都可见。这样既保证了存储器访问次序，而且又能确保更新结果都能影响到各个线程，每个核心的L1 Cache都会被更新。所以，使用原子操作做同步对于执行开销而言是相当大的，但是对于某些特定数据做同步更新时，不需要使用栅栏等这种更低效的处理机制，我们可以直接对那个存储地址采用原子操作。

OpenCL C 实现了C11的原子操作的子集，并且提供了非常丰富的原子操作种类，我们稍后会逐一详细讲解。不过，OpenCL 2.0之前的原子操作接口比较简单，而且与2.0版本完全不同，所以，我们这里先介绍一下OpenCL 1.2中的原子操作内建函数。

下面介绍一下OpenCL 1.2中的原子操作。

（1）原子递增

函数原型如下：
int atomic_inc (volitile _ _global int *p) unsigned int atomic_inc (volatile _ _global unsigned int *p)   int atomic_inc (volatile _ _local int *p) unsigned int atomic_inc (volatile _ _local unsigned int *p)

原子递增的操作数与返回类型是int或unsigned int,而可操作的存储空间可以是全局存储空间也可以是局部存储空间。此操作过程为：将参数p所指的地址内容取出，然后与1相加（即*p+1），最后将相加后的结果再写回p所指的地址中，然后返回原来修改前的p所指地址的内容。整个操作是原子的，即不可被打断的。

（2）原子递减

（3）原子加法

（4）原子减法 此三种方法的说明此处不再赘述。

七 OpenCL的地址空间

在OpenCL存储器模型中，我们知道OpenCL设备有全局存储器、局部存储器、常量存储器和私有存储器。对于这四种存储器，对应的地址空间修饰符为：_ global(或global)、 local(或local)、 constant(或constant)和 _private(或private)。如果一个变量由一个地址空间修饰符限定，那么这个变量就在指定的地址空间中分配。

程序中的函数参数，或者函数中缺省地址修饰符的局部变量，它们的地址修饰符为private。所有函数参数一定在私有地址空间。在程序范围内的一个变量，或者程序内的一个static变量，它们在全局或常量地址空间。如果没有地址修饰符制定，默认为全局的。

内核参数声明的指针类型必须指向global、local和constant三种类型之一。

内核函数返回类型必须是void类型，且只能在设备上执行。主机端可以调用这个函数。同时，如果一个内核函数调用另一个内核函数，那么被调的内核函数作为一个普通的函数调用。需要注意的是，如果内核函数中声明了local修饰符的变量，则在其他内核函数中调用此内核函数会有什么结果，这取决于OpenCL实现。

八 跋

由于OpenCL的细枝末节的地方较多，由于文章的篇幅有限，关于OpenCL简单介绍到这里，感兴趣的读者可以前往阅读刘文志《OpenCL异构并行计算》，同时知识星球「3Ｄ视觉技术交流群」也会不定期分享OpenCL的实践干货，欢迎加入一起学习～

欢迎加入公众号读者群，目前已有2D视觉、3D视觉、SLAM等微信群，请扫描下方微信二维码加群，备注：“昵称＋学校/公司＋研究方向”，，例如：“静静＋上海交大＋3D视觉”。请按如上格式备注，否则不予通过。

注：公众号「3D视觉工坊」后台回复「OpenCL」，即可获得关于OpenCL学习资料。

上述内容，如有侵犯版权，请联系作者，会自行删文。

星球成员，可免费提问，并邀进讨论群