Блог

Я пытаюсь написать ядро гистограммы в 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 memory. Пример гистограммы от nVidia разбивает рабочую нагрузку гистограммы на множество подзадач определенного размера, вычисляет их частичные гистограммы, а затем объединяет результаты в единую гистограмму. Не похоже, что это будет хорошо работать для изображений произвольного размера. Полагаю, я мог бы дополнить данные изображения фиктивными значениями…

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

Спасибо!

Как указывалось ранее, запись в общую память выполняется несинхронизированной и неатомной. Это приводит к ошибкам. Если изображение достаточно большое, у меня есть предложение:

Разделите свою рабочую группу на одномерную для столбцов или строк. Используйте каждое ядро для суммирования гистограммы для столбца или строки, а затем суммируйте ее глобально с помощью atom atom_inc. Это приводит к наибольшему количеству суммирования в частной памяти, что намного быстрее и сокращает атомарные операции.

Если вы работаете в двух измерениях, вы можете делать это по частям изображения.

[РЕДАКТИРОВАТЬ:]

Я думаю, у меня есть ответ получше : 😉

Взгляните на:http://developer.download.nvidia.com/compute/opencl/sdk/website/samples.html#oclHistogram

У них там интересная реализация…

Да, вы одновременно выполняете запись в общую память из многих рабочих элементов, поэтому вы потеряете элементы, если не выполните обновления безопасным способом (или хуже? Просто не делайте этого). Увеличение размера группы фактически увеличивает использование вашего вычислительного устройства, что, в свою очередь, увеличивает вероятность конфликтов. Таким образом, вы в конечном итоге теряете больше обновлений.

Однако вы, похоже, путаете синхронизацию (упорядочивание порядка выполнения потоков) и обновления совместно используемой памяти (которые обычно требуют либо атомарных операций, либо синхронизации кода и барьеров памяти, чтобы убедиться, что обновления памяти видны другим синхронизированным потокам).

барьер синхронизации не особенно полезен для вашего случая (и, как вы отметили, в любом случае недоступен для глобальной синхронизации. Причина в том, что 2 группы потоков могут никогда не запускаться одновременно, поэтому пытаться синхронизировать их бессмысленно). Обычно используется, когда все потоки начинают работать над созданием общего набора данных, а затем все начинают использовать этот набор данных с другим шаблоном доступа.

В вашем случае вы можете использовать атомарные операции (например, atom_inc, см. http://www.cmsoft.com.br/index.php?option=com_contentamp;view=categoryamp;layout=blogamp;id=113amp;Itemid=168). Однако обратите внимание, что обновление адреса памяти с высокой конкуренцией (скажем, из-за того, что у вас тысячи потоков, пытающихся выполнить запись всего в 256 целых), вероятно, приведет к снижению производительности. Все циклы, через которые проходит типичный код гистограммы, предназначены для уменьшения противоречий в данных гистограммы.

Вы можете проверить

• Пример гистограммы из AMD Accelerated Parallel Processing (APP) SDK.
• Глава 14 — Гистограмма изображения из книги «Руководство по программированию на OpenCL» (ISBN-13: 978-0-321-74964-2).

Гистограмма графического процессора — пример кода от Apple