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

最新推荐文章于 2025-12-01 22:16:11 发布

原创

最新推荐文章于 2025-12-01 22:16:11 发布 · 631 阅读

5 ·

CC 4.0 BY-SA版权

文章标签：

#ImageMagick #OpenCL #c++ #GraphicsMagick

作者分享了在实现ScaleImage的硬件加速过程中，通过理解ResizeHorizontalFilter函数和利用get_group_id来处理原始宽高的经验，展示了代码片段，描述了缩放图像时的局部内存管理和工作流设计，以及遇到的问题和解决方案。

文章目录

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

<2022-04-27 周三>

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

晚上做梦都在一直想这事儿，早上花了一个多小时，小有成果。

这里是参考ResizeHorizontalFilter()，居然把之前没有想明白的一些代码整清楚了：

accelerate.c:resizeHorizontalFilter()中传参gsize和lsize的地方，是拿目标宽高进行计算的，我的脑海中却一直用原始宽高去理解。

gsize[0]=(resizedColumns+pixelPerWorkgroup-1)/pixelPerWorkgroup*
  workgroupSize;
gsize[1]=resizedRows;
lsize[0]=workgroupSize;
lsize[1]=1;
outputReady=EnqueueOpenCLKernel(queue,horizontalKernel,2,
  (const size_t *) NULL,gsize,lsize,image,filteredImage,MagickFalse,
  exception);

这里的用意是在kernel函数ResizeHorizontalFilter()中可以方便的使用get_global_id(0)和get_local_id(0)，不用像我在“如何写ScaleImage()的硬件加速函数（四）”中那样去计算dst_local_w和filtered_x了。
但是似乎这里引出了另外一个问题，那原始宽高怎么用于计算呢？答案就是get_group_id(0)。

贴上目前开发一半的代码，效果是：

原图缩小一倍，水平方向显示原图全部，但被压缩一半；垂直方向显示原图上半部。
原图放大一倍，垂直方向显示原图全部，但被压缩一半；水平方向显示原图全部，但被拉长一倍。

附scaleFilter()和ScaleFilter()代码：

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,cl_uint rows,
  cl_mem scaledImageBuffer,cl_uint scaledColumns,cl_uint scaledRows,
  ExceptionInfo *exception)
{
   
   
  cl_kernel
    scaleKernel;

  cl_int
    status;

  const unsigned int
    workgroupSize = 256;

  float
    scale=1.0;

  int
    numCachedPixels;

  MagickBooleanType
    outputReady;

  size_t
    gammaAccumulatorLocalMemorySize,
    gsize[2],
    i,
    imageCacheLocalMemorySize,
    pixelAccumulatorLocalMemorySize,
    lsize[2],
    totalLocalMemorySize,
    weightAccumulatorLocalMemorySize;

  unsigned int
    chunkSize,
    pixelPerWorkgroup;

  int
    scale_ratio = 8; // related to the upper limit of zoom in?

  scaleKernel=NULL;
  outputReady=MagickFalse;

  scale=1.0/scale; // TODO(ocl)

  if (scaledColumns < workgroupSize)
  {
   
   
    chunkSize=32;
    pixelPerWorkgroup=32;
  }
  else
  {
   
   
    chunkSize=workgroupSize;
    pixelPerWorkgroup=workgroupSize;
  }

DisableMSCWarning(4127)
  while(1)
RestoreMSCWarning
  {
   
   
    /* calculate the local memory size needed per workgroup */
    numCachedPixels=(int) pixelPerWorkgroup;
    imageCacheLocalMemorySize=numCachedPixels*sizeof(CLQuantum)*4;
    totalLocalMemorySize=imageCacheLocalMemorySize;

    /* local size for the pixel accumulator */
    pixelAccumulatorLocalMemorySize=chunkSize*sizeof(cl_float4);
    totalLocalMemorySize+=pixelAccumulatorLocalMemorySize;

    /* local memory size for the weight accumulator */
    weightAccumulatorLocalMemorySize=chunkSize*sizeof(float);
    totalLocalMemorySize+=weightAccumulatorLocalMemorySize;

    /* local memory size for the gamma accumulator */
    gammaAccumulatorLocalMemorySize=chunkSize*sizeof