10.2 使用事件分析OpenCL代码

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

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

  1. cl_int clGetEventProfilingInfo(
  2.   cl_event event,
  3.   cl_profling_info param_name,
  4.   size_t param_value_size,
  5.   void *param_value,
  6.   size_t *param_value_size_ret)

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

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

事件状态 param_value返回的信息
CL_PROFILING_COMMAND_QUEUE 使用一个64位的值对主机端将命令提交到命令队列的时间进行统计(单位：ns)
CL_PROFILING_COMMAND_SUBMIT 使用一个64位的值对命令从命令队列提交到相关的设备上的时间进行统计(单位：ns)
CL_PROFILING_COMMAND_START 使用一个64位的值对命令开始的时间进行记录(单位：ns)
CL_PROFILING_COMMAND_END 使用一个64位的值对命令完成的时间进行记录(单位：ns)
CL_PROFILING_COMMAND_COMPLETE 使用一个64位的值对命令及其相关子命令完成的时间进行记录(单位：ns)

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

  1. // Sample code that can be used for timing kernel execution duration
  2. // Using different parameters for cl_profiling_info allows us to
  3. // measure the wait time
  4. cl_event timing_event;
  5. cl_int err_code;
  7. // !We ara timing the clEnqueueNDRangeKenrel call and timing
  8. // information will be stored in timing_event
  9. err_code = clEnqueueNDRangeKernel(
  10.   command_queue,
  11.   kernel,
  12.   work_dim,
  13.   global_work_offset,
  14.   global_work_size,
  15.   local_work_size,
  16.   0, NULL, &timing_event);
  18. cl_ulong starttime, endtime;
  20. err_code = clGetEventProfilingInfo(timing_event, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &starttime, NULL);
  21. kerneltimer = clGetEventProfilingInfo(timing_event, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &endtime, NULL);
  22. unsigned long elapsed = (unsigned long)(endtime - starttime);
  23. printf("Kernel Execution\t%ld ns\n", elapsed);

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