10.2 使用事件分析OpenCL代码

OpenCL命令队列支持64位的计时命令——使用clEnaueueXX()函数提交，比如:clEnqueueNDRangeKernel()。通常，命令入队都是异步的，并且开发者可以使用事件的方式对命令进行状态追踪。事件对象提供了一种方式来了解命令的执行过程。事件中记录了命令的很多相关信息，比如何时入队、何时提交到设备上、何时开始运行，以及何时执行完成。通过事件的信息获取函数——clGetEventProfilingInfo()，其能提供命令的相关计时信息：

使用事件对象显式的对OpenCL程序进行计，需要对对应的命令队列进行计时使能的操作。在创建命令队列的时候，需要设置CL_QUEUE_PROFILING_ENABLE标识。一旦命令命令队列创建完成，就无法在对事件计时的功能进行开启或关闭。

cl_int clGetEventProfilingInfo(
  cl_event event,
  cl_profling_info param_name,
  size_t param_value_size,
  void *param_value,
  size_t *param_value_size_ret)

第一个参数，event事件对象时必须给定的，第二个参数是一个枚举值，用来描述描述所要获取相应的时间信息。具体的值如表10.1所示。

表10.1 对应的命令状态可以用来获取OpenCL事件的时间戳

如之前所述，OpenCL命令队列是异步工作的，因此函数在命令入队时就返回了。所以在对事件对象进行计时查询时，需要调用一次clFinish()，以同步相关任务，让队列中的所有任务都完成。下面一段简单的代码展示了，如何使用事件的方式对内核进行性能分析。

{%ace edit=false, lang=’c_cpp’%} // Sample code that can be used for timing kernel execution duration // Using different parameters for cl_profiling_info allows us to // measure the wait time cl_event timing_event; cl_int err_code;

// !We ara timing the clEnqueueNDRangeKenrel call and timing // information will be stored in timing_event err_code = clEnqueueNDRangeKernel( command_queue, kernel, work_dim, global_work_offset, global_work_size, local_work_size, 0, NULL, &timing_event);

cl_ulong starttime, endtime;

err_code = clGetEventProfilingInfo(timing_event, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &starttime, NULL); kerneltimer = clGetEventProfilingInfo(timing_event, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &endtime, NULL); unsigned long elapsed = (unsigned long)(endtime - starttime); printf(“Kernel Execution\t%ld ns\n”, elapsed);

{%endace%}

代码清单10.1 使用OpenCL事件获取内核的时间信息