Каков наилучший способ обработки дополнительных данных, создаваемых небольшой частью потоков графического процессора в OpenCL?

Question

Я довольно новичок в OpenCL и имею следующую проблему:

У меня большой массив (6 * 1 000 000 поплавков). Для каждого элемента массива мне нужно сделать расчет. Базовый алгоритм хорошо работает на 16 графических процессорах (Tesla K80):

1.) Я создаю объект буфера массива и объект буфера для результатов для каждого устройства GPU и записываю его в каждую память GPU.

2.) Затем для каждого элемента массива создается поток, и вычисления выполняются в ядре на графических процессорах.

3.) Результат записывается в элемент массива результатов, соответствующий глобальному идентификатору потока.

4.) Хост читает буфер результатов.

Теперь я должен расширить этот алгоритм. Несколько элементов массива (10-100) действительно требуют дополнительного вычисления, которое дает дополнительный результат (еще 12 операций с плавающей запятой).

Вот некоторый псевдокод.

__kernel void calculation(__global float4 *input_array,
                          __global float4 *result_array){

    int id = get_global_id(0);

    //do calculation
    float4 result = some_func(input_array[id]);
    result_array[id] = result;


    if(some_rare_condition){
        //do another, much longer calculation
        float4 result2 = another_func(input_array[id]);
    }
}

Проблема в том, что у меня есть только несколько дополнительных результатов, и я не знаю, как лучше всего их сохранить и позволить хосту прочитать их.

Я не знаю, какие элементы массива требуют дополнительного вычисления, пока не вычислю первый результат.

Если бы это был C ++, я бы просто создал вектор для дополнительных результатов и вектор для индексов. Однако, насколько мне известно, в ядре OpenCL нет контейнеров динамической памяти.

Если я создаю второй массив результатов с 1 000 000 элементов и просто записываю в несколько требуемых позиций, это создает узкое место, когда я передаю его обратно хосту.

Если я создаю меньший массив, который определенно больше требуемого (например, 1000 элементов), я не уверен, как я могу позволить потокам писать в него безопасно.

Answer 1

Самое простое решение, вероятно, состоит в том, чтобы использовать атомный счетчик для назначения индексов в меньшем массиве.

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

Однако, действительно ли это эффективно , зависит от ряда факторов; что еще хуже, эти факторы работают друг против друга.

Прежде всего, звучит так, что вероятность того, что элемент из этого второго выходного массива будет меньше, составляет менее 0,1%. Небольшое число, подобное этому, полезно для атомных элементов: чем больше рабочих элементов хотят увеличить значение этого счетчика, тем выше вероятность того, что они попытаются сделать это одновременно и заблокировать друг друга и сериализовать. Кроме того, распределение имеет значение. Любая кластеризация будет работать против вас. Если ваши 0,1% в основном являются смежными, они также будут смежными в рабочей группе, а рабочие элементы в рабочей группе обычно выполняются в режиме блокировки на графическом процессоре. Поэтому им всем нужно подождать, пока друг друга завершат приращение, прежде чем любой из них сможет продолжить.

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

Возможные варианты простого подхода для устранения различных недостатков:

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

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