GraphicsMagick 的 OpenCL 开发记录（三十二）

本文详细解释了如何在OpenCL中编写ScaleImage的硬件加速函数，涉及clEnqueueNDRangeKernel函数的使用，以及如何结合__attribute__确保kernel函数调用。作者通过ResizeHorizontalFilter为例，剖析了ResizeImage内部逻辑并讨论了ScaleImage性能瓶颈及解决方案。

文章目录

如何写`ScaleImage()`的硬件加速函数（六）

<2022-04-27 周三>

如何写`ScaleImage()`的硬件加速函数（六）

不管什么事儿看来都怕琢磨，如果连做梦都能梦到你正在琢磨的事儿，估计离成功也就不远了。似乎目前已经达到了最好的效果，离目标越来越近了。

要理解clEnqueueNDRangeKernel()函数的第五第六个参数意义，但目前为止只能说暂时理解了。

cl_int clEnqueueNDRangeKernel (cl_command_queue command_queue,
  cl_kernel kernel,
  cl_uint work_dim,
  const size_t *global_work_offset,
  const size_t *global_work_size,
  const size_t *local_work_size,
  cl_uint num_events_in_wait_list,
  const cl_event *event_wait_list,
  cl_event *event)

第五第六个参数要结合__attribute__，否则无法调用kernel函数。

__kernel __attribute__((reqd_work_group_size(256, 1, 1)))

来回理解ResizeImage()和ScaleImage()函数的实现，对它们的内部逻辑流程了如指掌了可以说。
我选择以ResizeHorizontalFilter()为模板修改，之所以ResizeImage()处理速度慢，因为它的处理效果好，且有多种过滤效果可供选择，ResizeHorizontalFilter()是处理水平方向缩放，所以它将读入一整行原图像素，这正好和ScaleImage()的最内层循环处理方式相同。
ResizeHorizontalFilter()的最内层循环（如下），因为有累加操作，所以这正是处理水平缩放的操作。

for (unsigned int i = startStep; i < stopStep; i++, cacheIndex++)
{
   
   
  /* float weight = getResizeFilterWeight(resizeFilterCubicCoefficients,
    (ResizeWeightingFunctionType) resizeFilterType,
    (ResizeWeightingFunctionType) resizeWindowType,
    resizeFilterScale, resizeFilterWindowSupport,
    resizeFilterBlur, scale*(start + i - bisect + 0.5)); */

  float weight = getResizeFilterWeightForGM(resizeFilterType,
    scale*(start + i - bisect + 0.5), support);

  float4 cp = (float4)0.0f;

  __local CLQuantum *p = inputImageCache + (cacheIndex*4);
  cp.x = (float) *(p);
  cp.y = (float) *(p + 1);
  cp.z = (float) *(p + 2);

  if (matte_or_cmyk != 0)
  {
   
   
    cp.w = (float) *(p + 3);

    // float alpha = weight * QuantumScale * cp.w;
    // error: use of type 'double' requires cl_khr_fp64 support
    float alpha = weight * (1 - (float) cp.w / 255);

    filteredPixel.x += alpha * cp.x;
    filteredPixel.y += alpha * cp.y;
    filteredPixel.z += alpha * cp.z;
    filteredPixel.w += weight * cp.w;
    gamma += alpha;
  }
  else
    filteredPixel += ((float4) weight)*cp;

  density += weight;
}

“如何写ScaleImage()的硬件加速函数（五）”的问题在于没有办法处理图片下半部分（如何缩小一半的话），这里主要是因为y变量的限定（代码如下），因为传入kernel函数的gsize[1]=resizedRows;被限定的死死的。

const unsigned int pos = getPixelIndex(4, inputColumns, cacheRangeStartX, y);

目前只考虑按比例缩放，所以这里的y需要除以缩放比。
缩小后图片如果垂直方向相间着黑色宽竖条，那可能是因为numCachedPixels参数没有计算正确，这正可以修复“如何写ScaleImage()的硬件加速函数（五）”中的scale_ratio变量。

numCachedPixels=(int) ceil((pixelPerWorkgroup-1)/xFactor+2*support);

附上目前代码：

static MagickBooleanType scaleFilter(MagickCLDevice device,
  cl_command_queue queue,const Image *image,Image *filteredImage,
  cl_mem imageBuffer,cl_uint matte_or_cmyk,cl_uint columns