Как синхронизировать (конкретные) рабочие элементы на основе данных в OpenCL?

Question

Контекст:

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

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

Проблема:

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

Попытки:

В следующем примере каждый компонент может иметь максимум одинодин вход (для упрощения) превращает схему в дерево, где входом в схему является корень, а 3 выхода - это листы.

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

result сохраняет результат каждого компонента (напряжение, интенсивность и т. Д.)

inputModifiedуказать, если данный компонент уже рассчитал свой выходной сигнал.

// where the data come from (index in result)
constant int inputIndex[5]={0,0, 0, 2, 2};

kernel void update_component(
    local int *result, // each work-item result. 
    local int *inputModified // If all inputs are ready (one only for this example)
) {

    int id = get_local_id(0);
    int size = get_local_size(0);
    int barrierCount = 0;

    // inputModified is a boolean indicating if the input is ready
    inputModified[id]=(id!=0 ? 0 : 1);

    // make sure all input are false by default (except the first input).
    barrier(CLK_LOCAL_MEM_FENCE); 


    // Wait until all inputs are ready (only one in this example)
    while( !inputModified[inputIndex[id]] && size > barrierCount++)
    {
        // If the input is not ready, wait for it
        barrier(CLK_LOCAL_MEM_FENCE);
    }

    // all inputs are ready, compute output
    if (id!=0) result[id] = result[inputIndex[id]]+1;
    else result[0]=42;

    // make sure any other work-item depending on this is unblocked
    inputModified[id]=1;

    // Even if finished, we needs to "barrier" for other working items.
    while (size > barrierCount++)
    {
        barrier(CLK_LOCAL_MEM_FENCE);
    }
}

В этом примере N барьеров для N компонентов, что делает его хуже, чем последовательное решение.

Примечание: это только ядро,минимальный хост C ++ довольно длинный.В случае необходимости я мог бы найти способ добавить его.

Вопрос:

Возможно ли эффективно, и само ядро, чтобы различные рабочие элементы ожидаличтобы их данные были предоставлены другими рабочими элементами?Или какое решение будет эффективным?

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

Answer 1

Из документации барьера

https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/barrier.html

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

Но цикл while (содержащий барьер) в ядре имеет следующее условие:

inputModified[inputIndex[id]]

это может изменитьсяего поведение с идентификатором потока и приводит к неопределенному поведению.Кроме того, еще один барьер до этого

barrier(CLK_LOCAL_MEM_FENCE);

уже синхронизирует все рабочие элементы в рабочей группе, поэтому цикл while является избыточным, даже если он работает.

Также последний барьерный цикл

while (size > barrierCount++)
{
    barrier(CLK_LOCAL_MEM_FENCE);
}

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

Если вы хотите отправить какое-либо сообщение рабочим элементам вне рабочей группы, то вы можете использовать только атомарные переменные.Даже если вы используете атомикс, вы не должны принимать какой-либо порядок работы / выдачи между любыми двумя рабочими элементами.

Ваш вопрос

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

можно получить с помощью функции OpenCL 2.x «динамический параллелизм», которая позволяет рабочему элементупорождает новые рабочие группы / ядра внутри ядра.Это гораздо эффективнее, чем ожидание в цикле ожидания с вращением, и абсолютно более аппаратно-независимым, чем полагаться на количество потоков в полете, поддерживаемых графическим процессором (когда графический процессор не может обработать столько потоков в полете, любое ожидание вращения будетdead-lock, порядок потоков не имеет значения).

Когда вы используете барьер, вам не нужно сообщать другим потокам о "inputModified".Данные результата уже видны в рабочей группе.

Если вы не можете использовать OpenCL v2.x, вам следует обработать дерево с использованием BFS:

• начать 1 рабочий элемент для верхнего узла
• обрабатывает его, подготавливает K выходов и помещает их в очередь
• завершает ядро ​​
• запускает K рабочих элементов (каждый элемент pop из очереди)
• обрабатывает ихи подготовить N выходных данных и поместить их в очередь
• end kernel
• повторять до тех пор, пока в очереди больше не будет элементов

Количество вызовов ядра равно максимальномуглубина дерева, а не количество узлов.

Если вам нужна более быстрая синхронизация, чем «запуск ядра», то используйте одну рабочую группу для всего дерева, используйте барьер вместо вызовов ядра.Или обработайте первые несколько шагов на CPU, добавьте несколько поддеревьев и отправьте их в разные рабочие группы OpenCL.Возможно, вычисления на ЦП, пока есть N поддеревьев, где N = вычислительные единицы GPU, могут быть лучше для более быстрой асинхронной обработки поддеревьев на основе барьеров рабочей группы.

Существует также безбарьерная, безатомная и одиночнаяспособ вызова ядра для этого.Начните дерево снизу и поднимитесь.

Сопоставить все дочерние узлы самого глубокого уровня с рабочими элементами.Переместите каждого из них наверх во время записи их пути (идентификатор узла и т. Д.) В своей личной памяти / какой-либо другой быстрой памяти.Затем попросите их пройти сверху вниз по этому записанному пути, вычисляя на ходу, без каких-либо синхронизаций или даже атомик.Это менее эффективная работа, чем версии с барьером / вызовом ядра, но отсутствие барьера и наличие полностью асинхронных путей должно сделать его достаточно быстрым.

Если дерево имеет 10 глубин, это означает, что 10 указателей узлов для сохранения, а нетак много для частных регистров.Если глубина дерева составляет около 30 40, используйте локальную память с меньшим количеством потоков в каждой рабочей группе;если это еще больше, выделите глобальную память.

Но вам может потребоваться отсортировать рабочие элементы по топологии их пространственности / дерева, чтобы они работали вместе быстрее и с меньшим количеством ветвлений.

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

Если вам нужна только видимость данных для каждого рабочего элемента вместо группы или ядра, используйте fence: https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/mem_fence.html