OpenCL访问共享本地内存

Question

我编写了一个OpenCL内核程序，以对图像应用低通滤波器。内核：

__kernel void applyLowPassFilter(__global int *image, __global int *rst,
                                 __local int *localMem) {
  int nCols = get_global_size(0); // width of image
  int nRows = get_global_size(1); // height of image

  int xg = get_global_id(0); // x index of global buffer
  int yg = get_global_id(1); // y index od global buffer

  int xl = get_local_id(0); // x index of local buffer

  localMem[xl] = image[yg * nCols + xg];
  barrier(CLK_LOCAL_MEM_FENCE);
  if (yg != 0) {
    rst[yg * nCols + xg] = (localMem[xl] + image[(yg - 1) * nCols + xg]) / 2;
  }
}

在内核代码中，我想访问每个工作组的本地内存并计算值。因此，我将全局项目大小设置为W * H（W：图像的宽度，H：图像的高度），将本地项目大小设置为W * 1，我希望组大小为W，组数为大小在这里是H。主机代码：

    size_t globalItemSize[2];
    size_t localItemSize[2];
    globalItemSize[0] = W;
    globalItemSize[1] = H;
    localItemSize[0] = W;
    localItemSize[1] = 1;
    // Set cl kernel arguments.
    ret = clSetKernelArg(clKernel, 0, sizeof(cl_mem), (void *)&imageObj);
    ret = clSetKernelArg(clKernel, 1, sizeof(cl_mem), (void *)&rstObj);
    ret = clSetKernelArg(clKernel, 2, sizeof(int) * localItemSize[0], NULL); // local mem

但是，该代码无法正常工作，并且始终将结果图像显示为零。经过实验后，我发现它只能通过使用全局内存而不访问本地内存来工作。访问本地内存的代码我做错了吗？

Answer 1

我通过监视OpenCL return error codes知道了。首先，在CL_INVALID_KERNEL调用之后，我得到-48 clSetKernelArg错误代码，这表明我的内核出了点问题。然后，我删除了传递给访问本地内存的第三个参数，并在内核中使用了__local语句。此时，我收到了-51 CL_INVALID_ARG_SIZE错误代码，这使我想起了使用clinfo命令检查硬件的本地工作项数量限制的情况。意识到本地物品大小的限制，我将维度0中的localItemSize从W更改为W/3。

修改后的内核代码：

__kernel void applyLowPassFilter(__global int *image, __global int *rst) {
  int nCols = get_global_size(0); // width of image
  int nRows = get_global_size(1); // height of image

  int xg = get_global_id(0); // x index of global buffer
  int yg = get_global_id(1); // y index od global buffer

  int xl = get_local_id(0); // x index of local buffer

  __local int localMem[212]; // 1/3 of image width
  localMem[xl] = image[yg * nCols + xg];
  barrier(CLK_LOCAL_MEM_FENCE);
  if (yg != 0) {
    rst[yg * nCols + xg] = (localMem[xl] + image[(yg - 1) * nCols + xg]) / 2;
  }
}

主机代码中的参数配置：

    size_t globalItemSize[2];
    size_t localItemSize[2];
    globalItemSize[0] = W;
    globalItemSize[1] = H;
    localItemSize[0] = W / 3;
    localItemSize[1] = 1;
    // Set cl kernel arguments.
    ret = clSetKernelArg(clKernel, 0, sizeof(cl_mem), (void *)&imageObj);
    ret = clSetKernelArg(clKernel, 1, sizeof(cl_mem), (void *)&rstObj);