Я пытаюсь написать ядро гистограммы в OpenCL для вычисления 256 бинарных гистограмм 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 (1 877 400 пикселей), я склонен видеть только или около 1 870 000 записанных общих значений, когда суммирую значения гистограммы для каждого канала. Это также другое число каждый раз. Я ожидал этого, поскольку время от времени два ядра, вероятно, получают одно и то же значение из выходного массива и увеличивают его, фактически отменяя одну операцию увеличения (я предполагаю?).
Выход 1 870 000 предназначен для размера рабочей группы {1,1} (что, по-видимому, устанавливается по умолчанию, если я не укажу иное). Если я назначу больший размер рабочей группы, такой как {10,6}, я получу значительно меньшую сумму в моей гистограмме (пропорционально изменению размера рабочей группы). Мне это показалось странным, но я предполагаю, что происходит, когда все рабочие элементы в группе увеличивают значение выходного массива одновременно, и поэтому он просто считается одним приращением?
В любом случае, я читал в спецификации, что OpenCL не имеет глобальной синхронизации памяти, только синхронизацию внутри локальных рабочих групп, использующих их __local память. Пример гистограммы от nVidia разбивает рабочую нагрузку гистограммы на набор подзадач определенного размера, вычисляет их частичные гистограммы, а затем объединяет результаты в одну гистограмму. Это не похоже, что это будет хорошо работать для изображений произвольного размера. Я полагаю, я мог бы дополнить данные изображения фиктивными значениями ...
Будучи новичком в OpenCL, я думаю, мне интересно, есть ли более простой способ сделать это (так как кажется, что это должна быть относительно простая проблема GPGPU).
Спасибо!