Гистограмма изображения OpenCL

Я пытаюсь написать ядро ​​гистограммы в OpenCL для вычисления 256 bin R, G, и B-гистограммы входного изображения RGBA32F. Мое ядро ​​выглядит так:

const sampler_t mSampler = CLK_NORMALIZED_COORDS_FALSE |
                           CLK_ADDRESS_CLAMP|
                           CLK_FILTER_NEAREST;


__kernel void computeHistogram(read_only image2d_t input, __global int* rOutput,
                               __global int* gOutput, __global int* bOutput)
{

    int2 coords = {get_global_id(0), get_global_id(1)};

    float4 sample = read_imagef(input, mSampler, coords);

    uchar rbin = floor(sample.x * 255.0f);
    uchar gbin = floor(sample.y * 255.0f);
    uchar bbin = floor(sample.z * 255.0f);

    rOutput[rbin]++;
    gOutput[gbin]++;
    bOutput[bbin]++;


}

Когда я запускаю его на изображении 2100 x 894 (1877 400 пикселей), я, как правило, вижу только около 1870 000 общих записанных значений, когда я суммирую значения гистограммы для каждого канала. Также каждый раз это другое число. Я ожидал этого, поскольку время от времени два ядра, вероятно, захватывают одно и то же значение из выходного массива и увеличивают его, эффективно отменяя одну операцию приращения (я предполагаю?).

Вывод 1 870 000 предназначен для размера рабочей группы {1,1} (который, кажется, устанавливается по умолчанию, если я не укажу иное). Если я принудительно увеличиваю размер рабочей группы, например {10,6}, я получаю значительно меньшую сумму на моей гистограмме (пропорциональную изменению размера рабочей группы). Мне это показалось странным, но я Я догадываюсь, что происходит: все рабочие элементы в группе увеличивают значение выходного массива одновременно, и поэтому это считается одним приращением?

В любом случае, я читал в спецификации, что OpenCL не имеет синхронизация глобальной памяти, синхронизация только внутри локальных рабочих групп с использованием их __local памяти. Пример гистограммы от nVidia разбивает рабочую нагрузку гистограммы на группу подзадач определенного размера, вычисляет их частичные гистограммы, а затем объединяет результаты в одну гистограмму. Не похоже, что это будет хорошо работать с изображениями произвольного размера. Полагаю, я мог бы дополнить данные изображения фиктивными значениями ...

Поскольку я новичок в OpenCL, думаю, мне интересно, есть ли более простой способ сделать это (поскольку кажется, что это должен быть относительно простой GPGPU проблема).

Спасибо!