在英特尔核芯显卡与 C 主机代码上加速 OpenCL 代码的时间测量

Question

我是 openCL 的新手，愿意比较 C 代码和 openCL 内核之间的性能提升。 在将性能与 C 参考代码进行比较时，有人可以详细说明这两种方法中哪种方法更好/更正确地分析 openCL 代码：

1. 使用 QueryPerformanceCounter()/__rdtsc() 循环（在 getTime 函数中调用）

   ret |= clFinish(command_queue); //Empty the queue
   getTime(&begin);
   ret |= clEnqueueNDRangeKernel(command_queue, kernel, 2, NULL, global_ws, NULL, 0, NULL, NULL);  //Profiling Disabled.
   ret |= clFinish(command_queue);
   getTime(&end);
   g_NDRangePureExecTimeSec = elapsed_time(&begin, &end);      //Performs: (end-begin)/(CLOCK_PER_CYCLE*CLOCK_PER_CYCLE*CLOCK_PER_CYCLE)
   

2. 使用事件分析：

   ret = clEnqueueMarker(command_queue, &evt1);
   //Empty the Queue
   ret |= clEnqueueNDRangeKernel(command_queue, kernel, 2, NULL, global_ws, NULL, 0, NULL, &evt1);
   ret |= clWaitForEvents(1, &evt1);
   ret |= clGetEventProfilingInfo(evt1, CL_PROFILING_COMMAND_START, sizeof(cl_long), &begin, NULL);
   ret |= clGetEventProfilingInfo(evt1, CL_PROFILING_COMMAND_END, sizeof(cl_long), &end, NULL);
   g_NDRangePureExecTimeSec = (cl_double)(end - begin)/(CLOCK_PER_CYCLE*CLOCK_PER_CYCLE*CLOCK_PER_CYCLE);  //nSec to Sec
   ret |= clReleaseEvent(evt1);
   

此外，我没有使用专用显卡，而是将 Intel HD 4600 集成显卡用于以下 openCL 代码：

    __kernel void filter_rows(__global float *ip_img,\
                              __global float *op_img, \
                              int width, int height, \
                              int pitch,int N, \
                              __constant float *W)
    {
        __private int i=get_global_id(0); 
        __private int j=get_global_id(1); 
        __private int k;
        __private float a;
        __private int image_offset = N*pitch +N;
        __private int curr_pix = j*pitch + i +image_offset;

        // apply filter
        a  = ip_img[curr_pix-8] * W[0 ];    
        a += ip_img[curr_pix-7] * W[1 ];    
        a += ip_img[curr_pix-6] * W[2 ];    
        a += ip_img[curr_pix-5] * W[3 ];    
        a += ip_img[curr_pix-4] * W[4 ];    
        a += ip_img[curr_pix-3] * W[5 ];    
        a += ip_img[curr_pix-2] * W[6 ];    
        a += ip_img[curr_pix-1] * W[7 ];    
        a += ip_img[curr_pix-0] * W[8 ];    
        a += ip_img[curr_pix+1] * W[9 ];    
        a += ip_img[curr_pix+2] * W[10];    
        a += ip_img[curr_pix+3] * W[11];    
        a += ip_img[curr_pix+4] * W[12];    
        a += ip_img[curr_pix+5] * W[13];    
        a += ip_img[curr_pix+6] * W[14];    
        a += ip_img[curr_pix+7] * W[15];    
        a += ip_img[curr_pix+8] * W[16];
        // write output
        op_img[curr_pix] = (float)a;
    }

以及用于按列处理的类似代码。我使用方法 1 观察到增益（openCL Vs 优化矢量化 C-Ref）大约 11 倍，使用方法 2 大约 16 倍。 但是，我注意到有人声称使用专用显卡时的增益大约为 200-300 倍。

所以我的问题是：

1. 如果我在专用显卡上运行相同的代码，我可以获得多大的收益。会是类似的订单还是显卡会胜过英特尔高清显卡？
2. 我能否将 WARP 和线程概念从 CUDA 映射到英特尔高清显卡（即并行执行的线程数）？

Answer 1

您无法比较不同供应商的性能，基本的比较和预期可以通过并行线程运行的数量乘以它的频率来完成。

您有一个配备 Intel HD 4600 显卡的处理器：它应该有 20 个执行单元 (EU)，每个 EU 运行 7 个硬件线程，每个线程能够执行 SIMD8、SIMD16 或 SIMD32 指令，每个 SIMD 通道对应一个工作OpenCL 中的项目 (WI)。

SIMD16 是典型的简单内核，例如您要优化的内核，因此我们讨论的是 20*7*16=2240 个并行执行的工作项。请记住，每个工作项都能够处理矢量数据类型，例如float4，因此您绝对应该尝试重写内核以利用它们。我希望这也有助于您与 NVidia 的产品进行比较。

Answer 2

我观察到使用方法 1 的增益约为 11 倍，使用方法 2 时的增益约为 16 倍。

这看起来很可疑。在这两种情况下，您都使用高分辨率计数器。我认为您的输入尺寸太小并且会产生高运行变化。基于事件的测量稍微更准确，因为它在测量中不包括一些 OS + 应用程序开销。不过差别非常小。但在您的内核持续时间非常短的情况下，测量方法之间的差异......很重要。

如果我在专用显卡上运行相同的代码，我可以获得多大的收益。是类似订单还是显卡 会胜过英特尔高清显卡吗？

很大程度上取决于卡的功能。虽然英特尔核芯显卡是办公、电影和一些游戏的好显卡，但它无法与高端专用显卡相提并论。考虑到该卡具有非常高的功率包络、更大的芯片面积和更多的计算资源。预计专用卡将显示出更大的加速。 您的卡具有大约 600 GFLOPS 的峰值性能，而独立卡可以达到 3000 GFLOPS。 因此，您可以粗略地预计您的卡将比独立卡慢 5 倍。 但是，在说 300 倍加速时，请注意人们所比较的内容。如果他们与老一代 CPU 相比。他们可能是对的。但新一代 i7 CPU 确实可以缩小差距。

我可以将 WARP 和线程概念从 CUDA 映射到英特尔高清显卡（即并行执行的线程数）吗？

Intel HD 显卡没有扭曲。扭曲与 CUDA 硬件密切相关。基本上，warp 是相同的指令，由在 32 个 CUDA 核心上执行的 warp 调度程序调度。但是 OpenCL 与 CUDA 非常相似，因此您可以启动大量线程，这些线程将在您的显卡计算单元上并行执行。 但是在您的集成卡上编程时，最好忘记扭曲并知道您的卡有多少个计算单元。您的代码将在您的计算单元上并行运行在多个线程上。 换句话说，您的代码看起来与 CUDA 代码非常相似，但它将根据集成卡中可用的计算单元进行并行化。 然后每个计算单元可以以 SIMD 方式并行执行，例如.但 CUDA 的优化技术与编程 Intel HD 显卡的优化技术不同。