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 . Сначала я получил -48 CL_INVALID_KERNEL код ошибки после clSetKernelArg звонков. Очень подозрительно, что с моим ядром что-то не так. Затем я отбросил третий параметр, который я передал ядру для доступа к локальной памяти, и вместо этого использовал в коде ядра оператор __local. В этот момент я получил -51 CL_INVALID_ARG_SIZE код ошибки, который напомнил мне о проверке локального ограничения количества рабочих элементов моего оборудования с помощью команды clinfo. Поняв ограничение локального размера элемента, я изменил localItemSize в измерении 0 с 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);