与较大模板匹配的 OpenCL 模板比 OpenCV CPU 版本慢

Question

我是 opencl 的新手，现在我正在优化与 OpenCL 的模板匹配。我用较小的模板做了一些实验，发现我的 OpenCL 实现比 OpenCV 的 CPU 版本快。但在这种特殊情况下，模板尺寸非常大（2048x2048），原始图像尺寸为（3072x3072），OpenCV cpu 实现（137 秒）远远领先于 OpenCL（2000 秒）。请提出一些优化我的代码的方法，如下所示。

void __kernel corrln(global const unsigned char* ref_image, global const 
unsigned char* template, global float* corrln )
{
    const uint Width = get_global_size(0);
    const int2 pos = {get_global_id(0), get_global_id(1)};

    float sum = 0;

    for(int y = pos.y; y < 2048; y++ )
    {
       for(int x =pos.x; x < 2048; x++ )
       {
          const int2 xy = { x, y };
          const int2 txy = { x - pos.x, y - pos.y };
          sum += ref_image[index(xy, Width)] * template[index(txy, 
                 2048)];
      }
   }

  corrln[index(pos, Width)]= sum;               

}

Answer 1

考虑到您的 ref_image 的大小合理，小于 2048（例如 1024x1024），并且 ND 大小等于 ref_image 大小，每个 WI（工作项）都在进行不同数量的计算。

带有pos.x == 0 & pos.y == 0 的WI 在2 个循环内进行2048 * 2048 = 4M 计算，带有pos.x == 1023 & pos.y == 1023 的WI 在2 个循环内进行1024 * 1024M 计算。这对单身 WI 来说工作量太大了。

尝试以每个 WI 都会进行一些合理的固定数量计算的方式来简化这项任务。比如说，对于ref_image 的第一列，多次启动内核，每个内核将处理右侧的 16 列并计算和累加 corrln 数组，然后转到第二列，等等。

内核可能看起来像这样（仅用于说明！！！）：

void __kernel corrln(
    global const unsigned char* ref_image, 
    global const unsigned char* template, 
    global float* corrln ) 
{
    const uint Width = get_global_size(0);
    const int2 pos = {get_global_id(0), get_global_id(1)};
    uchar16 ref = vload16(index(xy, Width), ref_image);
    uchar16 tpl = vload16(index(xy, Width), template);
    float sum = corrln[index(pos, Width)] + dot(ref, tpl);
    corrln[index(pos, Width)]= sum;
}