测量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.

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);

3.等待内核完成

clWaitForEvents(1, &event);

4.等待所有排队的任务完成

clFinish(queue);

5.获取数据分析并计算内核执行时间(由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);

Answer 2

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

测量多次迭代总时间的最合适方法取决于“多次”在这里的含义。是NUMBER_OF_ITERATIONS=5还是NUMBER_OF_ITERATIONS=500000？如果迭代次数“很大”，您可以只使用系统时钟，可能使用特定于操作系统的功能，例如QueryPerformanceCounter在 Windows 上（另请参阅，例如，是否有一种方法可以使用 C 标准库测量高达微秒的时间？），但是当然，系统时钟的精度可能低于 OpenCL 设备的精度，因此这是否有意义实际上取决于迭代次数。

遗憾的是 NVIDIA 从他们的 Visual Profiler 中删除了 OpenCL 支持，不过……