OpenCL clock_gettime vs 内核分析：奇怪的结果

Posted 2023-03-06

【中文标题】OpenCL clock_gettime vs 内核分析：奇怪的结果

【英文标题】：OpenCL clock_gettime vs kernel profiling : strange results

【发布时间】：2014-06-03 10:27:32

【问题描述】：

我正在尝试分析简单卷积的不同实现。 我已经在不同的 CPU（i5、xeon 等）上获得了几个结果，现在我正在通过 intel beignet 尝试使用 intel HD4000。

我在主机端使用clock_gettime，在设备端使用CL_QUEUE_PROFILING_QUEUE 和事件。代码的精简版本是：

clock_gettime(CLOCK_REALTIME, &start);

err = clEnqueueNDRangeKernel(queue, img_conv_kernel, 2, NULL,
        &global_ws[0], &local_ws[0], 0, NULL, &event_clock);

if (err)
    die("can not launch kernel %d\n", err);

/* profiling */
clWaitForEvents(1, &event_clock);
clGetEventProfilingInfo(event_clock, CL_PROFILING_COMMAND_START, 
        sizeof(cl_ulong), &cl_start, NULL);
clGetEventProfilingInfo(event_clock, CL_PROFILING_COMMAND_END, 
        sizeof(cl_ulong), &cl_stop, NULL);

clock_gettime(CLOCK_REALTIME, &end);
printf("%f %f ", double) (cl_stop - cl_start) * 1e-6, 
            time_elapsed(start, end));

/* read data */
clock_gettime(CLOCK_REALTIME, &start);
err = clEnqueueReadBuffer(queue, res_d, CL_TRUE, 0, N*sizeof(float),
        res_h, 0, NULL, NULL);
clock_gettime(CLOCK_REALTIME, &end);

printf("%f ", time_elapsed(start, end));

/* C implementation */
clock_gettime(CLOCK_REALTIME, &start);
conv(img_data, res_h, &sobel_gx[0][0], k, k);
clock_gettime(CLOCK_REALTIME, &end);
printf("%f\n", time_elapsed(start, end));

结果是：

231.592960 16.701613 3.995006 151.874017
/* (device / host / reading-data / basic-c implementation )*/

我不明白的是，内核执行时间实际上大于通过clock_gettimecpu-time /em>，但根据 [0]，我使用 clWaitForEvents() 来确保内核已完全执行。

[0] : https://software.intel.com/en-us/articles/intel-sdk-for-opencl-applications-performance-debugging-intro

【问题讨论】：

能否请您澄清您到底在问什么并相应地编辑您的问题标题？

谢谢！我刚刚修好了。

您调用 printf() 2 次。首先输出 2 个值，然后输出 1 个值。你如何在标准输出中获得 4 个值？然后，OpenCL 分析结果以纳秒为单位，您将其转换为毫秒。你的clock_gettime的精度是多少，你是在比较毫秒和毫秒吗？

你在使用 OUT_OF_ORDER 队列吗？读取可能在内核完成之前完成。

@RomanArzumanyan 是的，为了便于阅读，我稍微修改了输出。是的，比较以毫秒为单位，我使用：((end.tv_sec - start.tv_sec) * 1e3 + (end.tv_nsec - start.tv_nsec) * 1e-6)

【参考方案1】：

请运行此代码并显示结果。

static long Time_Elapsed(
    long start,
    long end)

    return end - start;


static long Get_CL_Time(
    cl_event event)

    cl_ulong start, end;

    clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &start, NULL);     
    clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &end, NULL);

    return Time_Elapsed(start, end);


timespec start, end;
cl_event event_clock;

clock_gettime(CLOCK_REALTIME, &start);

/* run kernel */
err = clEnqueueNDRangeKernel(queue, img_conv_kernel, 2, NULL,
        &global_ws[0], &local_ws[0], 0, NULL, &event_clock);

clWaitForEvents(1, &event_clock);
long kernel_time = Get_CL_Time(event_clock);

/* read data */
err = clEnqueueReadBuffer(queue, res_d, CL_TRUE, 0, N*sizeof(float),
        res_h, 0, NULL, &event_clock);

clWaitForEvents(1, &event_clock);
long io_time = Get_CL_Time(event_clock);    

clock_gettime(CLOCK_REALTIME, &end);
long host_time = Time_Elapsed(start.tv_nsec, end.tv_nsec);

printf( "Kernel time: %l nanoseconds \n"
        "IO time:     %l nanoseconds \n"
        "Host time:   %l nanoseconds \n", 

        kernel_time,
        io_time,
        host_time);

/* C implementation */
clock_gettime(CLOCK_REALTIME, &start);
conv(img_data, res_h, &sobel_gx[0][0], k, k);
clock_gettime(CLOCK_REALTIME, &end);
host_time = Time_Elapsed(start.tv_nsec, end.tv_nsec);

printf("C implementation time: %l nanoseconds\n", host_time);

【讨论】：

这里是先生：Kernel time: 687377920 nanoseconds IO time: 53120 nanoseconds Host time: 53554859 nanoseconds C implementation time: 211760010 nanoseconds

您从设备读取的数据量是多少？内核执行时间和数据读取时间不同。

是的，这也是让我吃惊的地方。浮动中的高清图像，因此 1080*1920*sizeof(float)

因此，在这种情况下，您的下行速度应该是 ((1920 * 1080 * 4) * 1e9 / 53120) / (1024 * 1024 * 1024) = 145 Gb/s，速度快得令人难以置信。你确定内核没有产生错误吗？我的意思是不返回 NDRange 的状态，因为它是推入队列的状态。

刚刚检查过：我收到了一个缺少编译指示 (#pragma OPENCL EXTENSION cl_khr_fp64 : enable) 的警告，但仅此而已（它不会改变任何事情）。实际上我使用单色图像，所以它只有 1920*1080*1 但仍然......