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

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 与优化的矢量化 C-Ref 相比)约为 11 倍，使用方法 2 时的增益约为 16 倍。 不过，我注意到有人声称使用专用显卡时性能提升了 200-300 倍。

所以我的问题是:

1. 如果我在专用显卡中运行相同的代码，我可以获得多少增益。顺序会类似吗？或者显卡的性能会优于 Intel HD 显卡吗？
2. 我可以将 WARP 和线程概念从 CUDA 映射到英特尔高清显卡(即并行执行的线程数)吗？

Answer 1

I'm observing gain around 11x using method 1 and around 16x using method 2.

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

What magnitude of gain can I expect, if I run the same code in dedicated graphics card. Will it be similar order or graphics card will outperform Intel HD graphics?

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

Can i map WARP and thread concept from CUDA to Intel HD graphics (i.e. Number of threads executing in parallel)?

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