当前位置：极客笔记 > OpenCL 教程 > OpenCL 同步及事件机制 > OpenCL OpenCL事件机制

OpenCL OpenCL事件机制

OpenCL中，几乎所有clEnqueue系API包含了事件列表参数，可用于查询当前命令执行的状态以及对其后续命令的执行同步。

OpenCL提供了clCreateUserEvent函数API来创建用户自定义事件。如果使用clEnqueue将一个命令放入命令队列中，同时又把自己创建的用户自定义事件传进去，那么命令队列会等到此事件完成时（即CL_COMPLETE状态）才会将此命令提交给设备执行。当然，我们更多情况下是直接用clEnqueue系API所返回的cl_event对象来作为事件同步对象。

在我们用OpenCL事件对象来做命令执行的同步之前，我们先对如何查询OpenCL事件对象的状态以及相关注意事项做些了解。要查询当前OpenCL事件对象的状态，我们使用clGetEventInfo这个函数API，其原型是：

cl_int clGetEventInfo(cl_event         event,
                          cl_event_info   param_name,
                          size_t           param_value_size,
                          void ＊           param_value,
                          size_t ＊         param_value_size_ret);

这个函数返回对此函数调用的状态。如果是CL_SUCCESS，那么说明调用成功。如果是其他值，说明函数调用失败。
第1个参数event是指定要查询状态的OpenCL事件对象。第2个参数param_name指定了要查询事件的种类。就OpenCL 2.0而言，一共有以下几个事件种类：

  • CL_EVENT_COMMAND_QUEUE：表示查询此事件所对应的命令队列，因此第4个参数param_value对应于cl_command_queue类型。

  • CL_EVENT_CONTEXT：表示查询此事件所对应的OpenCL上下文，因此第4个参数param_value对应于cl_context类型。

  • CL_EVENT_COMMAND_TYPE：表示查询此事件所关联的命令类型，因此第4个参数param_value对应于cl_command_type类型。而cl_command_type类型所对应的值其实是枚举值，比如：CL_COMMAND_NDRANGE_KERNEL表示当前事件是与内核命令相关联的；CL_COMMAND_WRITE_BUFFER表示当前事件是与写缓存命令相关联的。由于命令种类繁多，因此这里不做详细介绍，各位读者可以参考OpenCL 2.0官方手册的5.11节。

  • CL_EVENT_COMMAND_EXECUTION_STATUS：表示查询此事件当前的执行状态，第4个参数param_value对应的类型为cl_int。可查询到的OpenCL命令的执行状态有以下4种：

    • CL_QUEUED：此状态说明当前命令已经被排入了命令队列中，但尚未提交。
    • CL_SUBMITTED：此状态说明当前命令已经被提交给了计算设备，准备执行。
    • CL_RUNNING：此状态说明当前计算设备正在执行此命令，但尚未完成。
    • CL_COMPLETE：此状态说明当前命令已经完成执行。

第3个参数param_value_size用于指定第4个参数param_value类型的大小。param_value即指向根据第2个参数param_name所指定的查询事件种类而要让clGetEventInfo函数最终输出结果的变量地址。第5个参数param_value_size_ret指向用于输出clGetEventInfo函数实际所返回的结果的大小，这个参数可以传空。下面我们举几个简单的例子来描述此API函数的用法。

【例1】 查询当前事件所对应的命令队列：

cl_command_queue queryCmdQueue;
clGetEventInfo(evt1, CL_EVENT_COMMAND_QUEUE,
                  sizeof(queryCmdQueue), &queryCmdQueue, NULL);

【例2】 查询当前事件所对应的OpenCL上下文：

cl_context queryContext;
clGetEventInfo(evt1, CL_EVENT_CONTEXT, sizeof(queryContext),
&queryContext, NULL);

下面我们将基于本章的样例代码，替换掉第1个clEnqueueWriteBuffer与第2个clEnqueueWriteBuffer调用之间的代码来描述查询使用非阻塞方式的clEnqueueWriteBuffer之后事件状态的变化。

//我们这里用evt1来监测对src1MemObj做数据传输的命令执行状态
cl_event evt1;
ret = clEnqueueWriteBuffer(command_queue, src1MemObj, CL_FALSE,
                                0, contentLength, pHostBuffer, 0, NULL,
                                &evt1);
if(ret != CL_SUCCESS)
{
    puts("Data1 transfer failed");
    goto FINISH;
}
cl_int status;
//我们用一个无限循环来观察事件在哪次迭代切换到了CL_SUBMITTED
for(int i = 0; ; i++)
{
    ret = clGetEventInfo(evt1, CL_EVENT_COMMAND_EXECUTION_STATUS,
                              sizeof(status), &status, NULL);
    if(ret == CL_SUCCESS)
    {
        if(status == CL_QUEUED)
        {
              printf("This write command has been queued: @%d\n", i);
              continue;
        }
        else if(status == CL_SUBMITTED)
        {
              printf("This write command has been submitted: @%d\n",
                      i);
              break;
        }
    }
}
clReleaseEvent(evt1);   //用完此事件对象后将其释放

上述代码中，我们在clEnqueueWriteBuffer函数的最后一个参数传入了在其上面声明的OpenCL事件对象的地址。一旦成功关联上，我们就获得了对应于当前写缓存命令的事件。然后，我们用一个无限循环来查询evt1事件的状态变化。这里要注意的是，对于不同系统环境下的OpenCL实现，事件查询返回的结果可能是不同的。在OS X 10.10下，第一次迭代中事件状态就已经是CL_QUEUED，说明此写命令已经被排入了命令队列里；过了几十次迭代后状态被切换为CL_SUBMITTED。而再继续，如果各位读者想要再增加对CL_RUNNING状态变化查询是无法得到结果的。因为CL_RUNNING状态与CL_COMPLETE状态其实不是在主机端维护的，而是在计算设备端，这一点与前两个状态不同。OS X中的OpenCL实现一般是在主机端上的驱动直接管理CL_QUEUED与CL_SUBMITTED状态，因为这两个状态都还没交付给计算设备。而后两种运行状态是计算设备处于此状态之后通过某种机制给主机端发送信号（如一个中断信号），然后主机端需要自己通过底层接口去获取此信号来更新状态。因此，在OS X的OpenCL实现中，如果要将当前的事件对象在计算设备端同步回来，必须调用clFinish或是后面将会介绍的clWaitForEvents这类同步API。
而在Windows与Linux下，如果使用AMD的OpenCL实现，在主机端可以捕获到CL_COMPLETE状态，而CL_RUNNING状态捕获不到。
另外比较重要的一点是，调用clEnqueueWriteBuffer函数所在的线程必须与调用clGetEventInfo函数所在的线程是同一个。如果clEnqueueWriteBuffer函数在主线程上调用，而clGetEventInfo函数在另一个用户线程上调用，则执行clGetEventInfo函数时可能会引发异常。
上述示例代码中，我们使用了一个无限循环来轮询当前事件的状态。OpenCL还提供了一种异步回调的方式来跟踪当前事件的状态变化。其原型是：

cl_int clSetEventCallback(cl_event    event,
                              cl_int      command_exec_callback_type,
                        void (＊pfn_notify)(cl_event, cl_int, void ＊),
                              void ＊      user_data);

其第一个参数event为指定的事件对象。第二个参数command_exec_callback_type用来指明当此事件处于哪个执行状态时发生回调。可注册的执行状态有CL_SUBMITTED、CL_RUNNING和CL_COMPLETE。第三个参数pfn_notify就是我们自己定义的观察事件状态的回调函数。其三个参数分别指定OpenCL系统传入的事件对象，此事件当前的状态以及用户数据参数。这个用户数据参数与clSetEventCallback函数的第四个参数user_data是同一个值。
下面我们对上述代码例子再做些小修改来测试这个API函数。首先，我们在main函数的上面定义以下内容：

static volatile bool canContinue = false;
static void MyEventHandler(cl_event event, cl_int status,
                                void ＊userData)
{
}
if(status == CL_SUBMITTED)
puts("The current status is submitted.");
canContinue = true;

上述的canContinue全局变量将会用于判定程序是否可以继续往下执行。然后，我们用下面的代码来替换上面for循环部分的代码：

clSetEventCallback(evt1, CL_SUBMITTED, &MyEventHandler, NULL);
for(int i = 0; ; i++)
{
}
if(canContinue)
{
printf("This is the %dth iteration.\n", i);
break;
}

我们通过clSetEventCallback函数来监测evt1事件对象状态是否切换到了CL_SUBMITTED。当evt1事件对象的状态变为CL_SUBMITTED，那么OpenCL系统将会调用我们自定义的MyEventHandler函数。我们在MyEventHandler函数中将canContinue全局标志设置为true，用于指明for循环的轮询可以跳出。
下面我们来讲述如何通过事件对象做命令之间的同步。大部分clEnqueue系函数API都含有指定等待事件列表的参数，也就是const cl_event *event_wait_list。如果此参数不空，并且所指向的存储空间包含了N个事件对象（N 〉 0），那么该enqueue命令会在此等待事件列表中的所有事件都处于CL_COMPLETE之后才会提交给计算设备执行。
下面，我们将基于本章节的示例代码，修改从第一个clEnqueueWriteBuffer函数调用到FINISH跳转标签上面的所有代码，替换如下：

//我们这里用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 = clEnqueueWriteBuffer(command_queue, src2MemObj, CL_TRUE, 0,
                                contentLength, pHostBuffer, 1, &evt1,
                                &evt2);
if(ret != CL_SUCCESS)
{
    puts("Data2 transfer failed");
    goto FINISH;
}
clReleaseEvent(evt1);
clReleaseEvent(evt2);

我们先用evt1事件对象作为第一个写缓存命令的同步事件。然后，在调用第二个clEnqueueWriteBuffer函数时把evt1对象的地址作为等待事件列表的参数，而evt2作为第二个写缓存命令的同步事件对象。这样就使得在第二个写缓存命令开始执行前，第一个写缓存命令必须完成执行。
上面介绍了命令之间的同步。那么如果我们的需求是在某个命令完成之前不想让当前主机端的线程继续往下执行该怎么办呢？如果使用clFinish函数，那么主机端的线程会被一直挂起，直到命令队列中所有命令全都执行完了之后才能返回操作。而如果我们仅仅只是等某一个命令执行完成，就可以使用clWaitForEvents函数接口，其声明如下：

cl_int clWaitForEvents (cl_uint num_events,
                            const cl_event ＊event_list)

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

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: %ldus\n", duration);

我们先用阻塞方式调用第2个clEnqueueWriteBuffer函数，察看结果；然后再用非阻塞方式调用第2个clEnqueueWriteBuffer函数，察看结果。在本章使用示例代码的环境下，使用阻塞方式仅等待了4微秒，说明clWaitForEvents的函数执行很快就被返回了，因为阻塞方式下，evt2事件在调用clEnqueueWriteBuffer函数时就已经处于完成状态了；而使用非阻塞的方式则等待了36 601微秒，说明clWaitForEvents函数的执行将当前线程挂起了相当一段时间等待evt2处于CL_COMPLETE状态。
正如上面所提到的，如果我们在clWaitForEvents的函数下面添加clGetEventInfo函数调用来查询当前事件的执行状态，我们会看到，evt2一定处于CL_COMPLETE状态（被定义为0）：

cl_int status;
ret = clGetEventInfo(evt2, CL_EVENT_COMMAND_EXECUTION_STATUS,
                        sizeof(status), &status, NULL);
printf("The current status of evt2 is: %d\n", status);

了解OpenCL几个典型的同步用法之后，下面我们将举一个更综合性的例子将这些同步函数API的使用整合在一起。我们把本章节一开始提供的示例代码从struct timeval tsBegin, tsEnd；一直到FINISH跳转标签之间的代码替换为以下代码，同时将evt1和evt2的声明放到了前面，与其他变量写在一起，各位读者在FINISH标签下要增加对这两个事件对象的释放：

/＊对src1MemObj的数据传输,我们使用非阻塞方式,
等后续设置完成后通过事件等待机制进行同步＊/
ret = clEnqueueWriteBuffer(command_queue, src1MemObj, CL_FALSE,
                                0, contentLength, pHostBuffer, 0, NULL,
                                &evt1);
if(ret != CL_SUCCESS)
{
    puts("Data1 transfer failed");
    goto FINISH;
}
/＊ 对src2MemObj的数据传输,我们使用非阻塞方式,
等后续设置完成后通过事件等待机制进行同步＊/
ret = clEnqueueWriteBuffer(command_queue, src2MemObj, CL_FALSE,
                                0, contentLength, pHostBuffer, 1, &evt1,
                                &evt2);
if(ret != CL_SUCCESS)
{
    puts("Data2 transfer failed");
    goto FINISH;
}
//创建用于结果输出的缓存对象
//我们这里使用可读可写是为了在第一个kernel程序执行完之后,
//它既能作为第二个kernel程序的输入,也能作为第二个kernel程序的输出
dstMemObj = clCreateBuffer(context, CL_MEM_READ_WRITE,
                                contentLength, NULL, &ret);
//指定内核源文件路径,这个路径根据读者当前环境可以更改
//这里使用绝对路径也是避免不同系统需要调用不同API来获取当前路径
const char ＊pFileName = "/Users/zennychen/Downloads/test.cl";
FILE ＊fp = fopen(pFileName, "r");
if (fp == NULL)
{
    puts("The specified kernel source file cannot be opened!");
    goto FINISH;
}
fseek(fp, 0, SEEK_END);
const long kernelLength = ftell(fp);
fseek(fp, 0, SEEK_SET);
kernelSource = malloc(kernelLength);
fread(kernelSource, 1, kernelLength, fp);
fclose(fp);
program = clCreateProgramWithSource(context, 1,
                                          (const char ＊＊)&kernelSource,
                                          (const size_t ＊)&kernelLength,
&ret);
ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);
if (ret != CL_SUCCESS)
{
    size_t len;
    char buffer[8 ＊ 1024];
    printf("Error: Failed to build program executable!\n");
    clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG,
                              sizeof(buffer), buffer, &len);
    printf("%s\n", buffer);
    goto FINISH;
}
//第一个kernel的主函数为kernel1_test
kernel = clCreateKernel(program, "kernel1_test", &ret);
if(kernel == NULL)
{
    puts("Kernel failed to create!");
    goto FINISH;
}
ret = clSetKernelArg(kernel, 0, sizeof(cl_mem),
                        (void ＊)&dstMemObj);
ret |= clSetKernelArg(kernel, 1, sizeof(cl_mem),
                          (void ＊)&src1MemObj);
ret |= clSetKernelArg(kernel, 2, sizeof(cl_mem),
                          (void ＊)&src2MemObj);
if(ret != CL_SUCCESS)
{
    puts("Set arguments error!");
    goto FINISH;
}
//获取最大工作组大小
size_t maxWorkGroupSize = 0;
clGetDeviceInfo(device_id, CL_DEVICE_MAX_WORK_GROUP_SIZE,
                  sizeof(maxWorkGroupSize), &maxWorkGroupSize, NULL);
/＊我们这里等待对src1MemObj和src2MemObj的数据
全都传输好之后再执行下面的内核程序执行＊/
clWaitForEvents(2, (cl_event[2]){evt1, evt2});
//这里用完evt1与evt2之后将它们释放置空
clReleaseEvent(evt1);
clReleaseEvent(evt2);
evt1 = NULL;
evt2 = NULL;
//这里指定将总共有(contentLength / sizeof(int))个工作项
//然后,每个工作组含有maxWorkGroupSize个工作项
//我们这里再复用evt1来跟踪内核程序1的执行状态
ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL,
                                  (const size_t[])
                                  {contentLength / sizeof(int)},
                                  &maxWorkGroupSize, 0, NULL, &evt1);
if(ret != CL_SUCCESS)
{
}
    puts("kernel1 execution failed");
    goto FINISH;
//下面,我们初始化第二个kernel程序
//现在计算设备在做计算时,我们主机端能不受干扰地继续做其他事情
kernel2 = clCreateKernel(program, "kernel2_test", &ret);
//设置kernel2_test的参数
ret = clSetKernelArg(kernel2, 0, sizeof(cl_mem),
                        (void ＊)&dstMemObj);
ret |= clSetKernelArg(kernel2, 1, sizeof(cl_mem),
                          (void ＊)&src1MemObj);
ret |= clSetKernelArg(kernel2, 2, sizeof(cl_mem),
                          (void ＊)&src2MemObj);
if(ret != CL_SUCCESS)
{
    puts("Kernel2 arguments setting failed");
goto FINISH;
}
//这里kernel2程序必须等kernel1执行完成之后才能执行
ret = clEnqueueNDRangeKernel(command_queue, kernel2, 1, NULL,
                              (const size_t[])
                              {contentLength / sizeof(int)},
                              &maxWorkGroupSize, 1, &evt1, &evt2);
if(ret != CL_SUCCESS)
{
puts("kernel2 execution failed");
goto FINISH;
}
//准备做校验
pDeviceBuffer = (int ＊)malloc(contentLength);
//这里,读取计算设备端的数据的命令通过evt2进行同步
//确保kernel2完成执行后再执行读数据命令,并且这里使用阻塞的方式读取数据
clEnqueueReadBuffer(command_queue, dstMemObj, CL_TRUE, 0,
                  contentLength, pDeviceBuffer, 1, &evt2, NULL);
for(int i = 0; i 〈 contentLength / sizeof(int); i++)
{
int testData = pHostBuffer[i] + pHostBuffer[i];
testData = testData ＊ pHostBuffer[i] - pHostBuffer[i];
if(testData != pDeviceBuffer[i])
{
    printf("Error occurred @%d, result is: %d\n", i,
            pDeviceBuffer[i]);
    goto FINISH;
}
}
puts("Result is OK!");

上述代码中我们结合了clEnqueue自身的事件同步机制，读写数据的阻塞与非阻塞方式，以及事件等待函数的调用来做各种情况下的同步。下面列出上述函数中所用到的内核代码，也非常简单：

__kernel void kernel1_test(__global int ＊pDst,
                                __global int ＊pSrc1, __global int ＊pSrc2)
{
    int index = get_global_id(0);
    pDst[index] = pSrc1[index] + pSrc2[index];
}
__kernel void kernel2_test(__global int ＊pDst,
                                __global int ＊pSrc1, __global int ＊pSrc2)
{
    int index = get_global_id(0);
    pDst[index] = pDst[index] ＊ pSrc1[index] - pSrc2[index];
}

本节给大家讲述了OpenCL命令之间的同步以及命令与主机端线程同步的用法和实例。

Camera课程

Camera API2 教程

API2 教程
Camera 性能教程

Perf 教程
Camera Native Framework 教程

Native FW 教程
Camera QCOM ISP 教程

ISP 教程
Camera QCOM Sensor 教程

Sensor 教程

Python教程

Python 教程

Python 教程
NumPy 教程

NumPy 教程
Django 教程

Django 教程

Java教程

Java 教程

Java 教程

Web教程

JavaScript 教程

JavaScript 教程
jQuery 教程

jQuery 教程
CSS 教程

CSS 教程
PHP 教程

PHP 教程
Laravel 教程

Laravel 教程
TypeScript 教程

TypeScript 教程
WordPress 教程

WordPress 教程

数据库教程

SQL 教程

SQL 教程
SQLite 教程

SQLite 教程
MySQL 教程

MySQL 教程
PL/SQL 教程

PL/SQL 教程
MongoDB 教程

MongoDB 教程

图形图像教程

OpenCV 教程

OpenCV 教程

办公软件教程

Excel 教程

Excel 教程

Linux教程

Linux命令全集

Linux命令全集
Linux内核API

Linux内核API
LVGL 教程

LVGL 教程

计算机教程

操作系统 教程

操作系统 教程
计算机网络 教程

计算机网络 教程
C语言 教程

C语言 教程
C++ 教程

C++ 教程
Swift 教程

Swift 教程

大数据教程

Hadoop 教程

Hadoop 教程
Spark 教程

Spark 教程
Scala 教程

Scala 教程

开发工具教程

Git 教程

Git 教程
Jenkins 教程

Jenkins 教程
ChatGPT 教程

ChatGPT 教程
IntelliJ 教程

IntelliJ 教程

OpenCL 同步及事件机制