测量OpenCL内核的执行时间

Question

我有下面的循环来测量我的内核的时间：

double elapsed = 0;
cl_ulong time_start, time_end;
for (unsigned i = 0; i < NUMBER_OF_ITERATIONS; ++i)
{
    err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0, NULL, &event); checkErr(err, "Kernel run");
    err = clWaitForEvents(1, &event); checkErr(err, "Kernel run wait fro event");
    err = clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_START, sizeof(time_start), &time_start, NULL); checkErr(err, "Kernel run get time start");
    err = clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_END, sizeof(time_end), &time_end, NULL); checkErr(err, "Kernel run get time end");
    elapsed += (time_end - time_start);
}

然后我将 elapsed 除以 NUMBER_OF_ITERATIONS 得到最终估计值。但是，我担心单个内核的执行时间太短，可能会对我的测量结果产生不确定性。如何测量所有 NUMBER_OF_ITERATIONS 内核组合所需的时间？

您能否推荐一种分析工具来帮助解决这个问题，因为我不需要通过编程方式访问这些数据。我使用 NVIDIA 的 OpenCL。

- user1096294

Answer 1

您需要按照以下步骤来测量OpenCL内核执行时间：

1. 创建一个队列，在创建队列时启用性能分析：

cl_command_queue command_queue;
command_queue = clCreateCommandQueue(context, devices[deviceUsed], CL_QUEUE_PROFILING_ENABLE, &err);

2. 在启动内核时链接一个事件

cl_event event;
err=clEnqueueNDRangeKernel(queue, kernel, woridim, NULL, workgroupsize, NULL, 0, NULL, &event);

等待内核执行完毕。

clWaitForEvents(1, &event);

3. 等待所有进入队列的任务完成

clFinish(queue);

4. 获取分析数据并计算内核执行时间（由OpenCL API以纳秒返回）

cl_ulong time_start;
cl_ulong time_end;

clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_START, sizeof(time_start), &time_start, NULL);
clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_END, sizeof(time_end), &time_end, NULL);

double nanoSeconds = time_end-time_start;
printf("OpenCl Execution time is: %0.3f milliseconds \n",nanoSeconds / 1000000.0);

- Dongwei Wang

Answer 2

分析函数返回纳秒，非常精确（约为50纳秒），但是执行时间取决于其他您无法控制的小问题而不同。

这减少了您希望测量的问题：

• 测量内核执行时间：您的方法是正确的，平均执行时间的精度随着N的增加而增加。仅考虑执行时间，不考虑开销。

• 测量内核执行时间+开销：您也应使用事件，但应自CL_PROFILING_COMMAND_SUBMIT以来进行测量，以考虑额外的执行开销。

• 测量实际主机端执行时间：您也应使用事件，但应自第一个事件开始测量，直到最后一个事件结束。使用CPU计时测量是另一种可能性。如果要进行此类测量，则应从循环中删除等待事件，以允许OpenCL系统的最大吞吐量（并尽可能减少开销）。

回答工具问题，我建议使用nVIDIA可视化分析器。但由于它不再适用于OpenCL，您应该使用Visual Studio插件或nvprofiler的旧版本（CUDA 3.0）。

- DarkZeros

Answer 3

被测量的时间以纳秒为单位返回，但你说得对：计时器的分辨率较低。不过，当你说时间太短无法准确测量时，我想知道你的内核实际执行时间是多少（我的直觉是分辨率应该在微秒级别范围内）。

最适合测量多次迭代总时间的方法取决于这里的“多次”是什么意思。是NUMBER_OF_ITERATIONS=5还是NUMBER_OF_ITERATIONS=500000？如果迭代次数“很大”，你可以使用系统时钟，可能需要使用特定于操作系统的函数，例如Windows上的QueryPerformanceCounter（也可以参见Is there a way to measure time up to micro seconds using C standard library?），但当然，系统时钟的精度可能低于OpenCL设备，因此这是否有意义取决于迭代次数。

不过，遗憾的是NVIDIA从其Visual Profiler中删除了OpenCL支持...

- Marco13

---

谢谢你的回答！我们正在讨论数十微秒，而NUMBER_OF_ITERATIONS实际上是任意的。目前我使用的是30。我需要一个好的数字来消除错误。 - user1096294

Answer 4

在英特尔的OpenCL GPU实现中，我已经成功地使用了您的方法（每个内核的计时），并且更喜欢它而不是批处理NDRanges的流。

• 另一种方法是使用标记事件运行N次，并测量时间，就像this question（问题而不是答案）中提出的方法。

• 根据我的经验，短内核的时间通常至少在微秒级别。

• 您可以使用CL_DEVICE_PROFILING_TIMER_RESOLUTION（例如，在我的设置中为80 ns）使用clGetDeviceInfo检查计时器分辨率。

- Tim