关于OpenCL总执行时间过长及与设备时间不匹配的技术问询

打印OpenCL事件各阶段时间的函数 #

void PrintProfilingInfo(cl_event event)
{
    cl_int err_num = -1;

    cl_ulong t_queued;
    cl_ulong t_submitted;
    cl_ulong t_started;
    cl_ulong t_ended;
    cl_ulong t_completed;

    err_num = clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_QUEUED,
                                      sizeof(cl_ulong), &t_queued, NULL);

    // submit time
    err_num = clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_SUBMIT,
                                      sizeof(cl_ulong), &t_submitted, NULL);

    // start time
    err_num = clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_START,
                                      sizeof(cl_ulong), &t_started, NULL);

    // end time
    err_num = clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_END,
                                      sizeof(cl_ulong), &t_ended, NULL);

    // complete time
    err_num = clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_COMPLETE,
                                      sizeof(cl_ulong), &t_completed, NULL);

    printf("queue  -> submit : %fus\n", (t_submitted - t_queued) * 1e-3);
    printf("submit -> start  : %fus\n", (t_started - t_submitted) * 1e-3);
    printf("start  -> end    : %fus\n", (t_ended - t_started) * 1e-3);
    printf("end    -> finish : %fus\n", (t_completed - t_ended) * 1e-3);
}

计算总执行时间的代码片段 #

timeval t_start;
long long time_diff;
timeval end;

gettimeofday(&t_start, NULL);

err_code = clEnqueueNDRangeKernel(cl_cmd_queue,
                                    kernel,
                                    2,
                                    NULL,
                                    global_work_size,
                                    local_work_size,
                                    0,
                                    NULL,
                                    &kernel_event);
err_code = clWaitForEvents(1, &kernel_event);

gettimeofday(&end, NULL);
time_diff = 1000000 * (end.tv_sec - t_start.tv_sec) + end.tv_usec - t_start.tv_usec;
printf(" ==>> OpenCL Gaussian Blur average cost: %lld us\n", time_diff);

（注：原代码中start.tv_sec存在笔误，已修正为t_start.tv_sec）

执行结果截图 #

技术疑问 #

1. 为何总执行时间耗时过长？
2. 为何设备各阶段时间总和与总执行时间不相等？队列到提交阶段的主要耗时环节是什么？

已查阅OpenCL官方文档，未找到队列到提交阶段的详细说明。

问题解答 #

1. 总执行时间过长的可能原因

• 主机-设备数据传输开销：如果高斯模糊涉及大量主机与设备间的数据拷贝（如clEnqueueReadBuffer/clEnqueueWriteBuffer），这部分操作未被事件profiling统计，但会被gettimeofday的总计时覆盖，是总时间过长的常见原因。
• 队列阻塞或同步等待：若命令队列是阻塞模式（未开启CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE），或之前有未完成的命令占用设备资源，clEnqueueNDRangeKernel或clWaitForEvents会等待资源释放，这部分等待时间会被计入总执行时间。
• 设备初始化/资源分配延迟：如果是程序首次执行OpenCL kernel，设备需完成上下文初始化、kernel即时编译、工作组内存分配等操作，这些开销由主机计时承担，但不会被事件阶段时间统计到。
• 系统资源竞争：主机CPU同时处理其他任务，或GPU被其他进程占用，都会拉长总执行时间，而设备侧profiling仅统计自身执行环节的耗时。

2. 设备阶段时间总和与总执行时间不相等的原因及队列→提交阶段的耗时环节

时间不相等的核心原因

主机侧gettimeofday计时覆盖主机端所有操作开销，而设备事件profiling仅统计设备内部的阶段时间，统计范围完全不同：

• 主机计时包含：kernel入队前的参数检查、主机端命令队列调度开销、等待设备完成命令的主机线程阻塞时间、甚至主机操作系统的调度延迟。
• 设备事件时间戳是设备硬件/驱动层面记录的，仅追踪命令在设备侧的流转，不包含主机端任何操作耗时。

队列（QUEUED）→提交（SUBMIT）阶段的主要耗时环节

这个阶段是主机端将命令从用户空间提交到设备驱动/硬件队列的过程，主要耗时点包括：

• 主机端命令验证：驱动检查kernel参数合法性、工作组大小是否符合设备限制、内存对象是否有效等。
• 命令队列同步与调度：若为顺序执行队列，驱动需等待队列中之前的命令完成（或达到同步点），才能将当前命令提交给设备。
• 主机到设备的命令传输：将kernel配置参数、工作尺寸等信息从主机内存拷贝到设备命令缓冲区。
• 驱动内部资源准备：比如为kernel分配临时设备内存、绑定相关内存对象到硬件单元等。

OpenCL官方文档对该阶段描述模糊，是因为不同厂商驱动实现差异极大，具体耗时环节依赖硬件架构和驱动优化程度。

内容的提问来源于stack exchange，提问作者lanling