OpenCL, выполнение Синхронизация WorkItems в рабочей группе (не барьеров и обращение к памяти)

Question

@ Привет всем,

Я реализует цифровой фильтр на AMD GPU с помощью OpenCL. Функция фильтра - это зависимости между соседними элементами. Каждый элемент зависит от элементов слева, сверху и сверху справа. См. LINK с изображением, чтобы получить его лучше.

Такие,

element 2 -> from 1. 
element 3 -> from 2. 
element 7 -> from 1, 2 (top, top-right). 
element 8 -> from 7, 2, 3. 
element 16 -> from 15, 10, 11 
etc... applying for each element from 2nd to 24th 

Всего в нескольких подряд может выполняться параллельно, если выше строка идти вперед на 2-х элементов.

Итак, я пришел к вопросу, как организовать синхронизацию между элементами ??? Я не нашел способов организовать синхронизацию внутри рабочей группы. Барьеры, кажется, не являются подходящим механизмом, поскольку они предназначены для одновременного хранения контрольных точек архивирования, в то время как в моем случае это должна быть синхронизация с задержкой между потоками.

Существует один способ, который кажется подходящим. Использование clEnqueueTask() и событий (или clEnqueueNDRangeKernel с рабочей группой равно < 1, 1, 1>). Но в этом случае I не может использовать локальную память, и это будет значительно медленнее.

Вопрос еще раз: как организовать синхронизацию между потоками в одной рабочей группе для эффективного использования быстрой локальной памяти ???

Заранее спасибо.
Обратите внимание:
- размер 6x4 только для иллюстрации, действительно, это около 4K элементов.
- "element" - это блок 8x8 целых значений.

Answer 1

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

element 2 -> from 1. 
BARRIER -------------------------------------------(SYNC WG) 
element 3 -> from 2. 
BARRIER -------------------------------------------(SYNC WG) 
element 7 -> from 1, 2 (top, top-right). 
BARRIER -------------------------------------------(SYNC WG) 
element 8 -> from 7, 2, 3. 
BARRIER -------------------------------------------(SYNC WG) 
element 16 -> from 15, 10, 11. 

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

.____WG1_____.____WG2_____.____WG3_____.____WG4_____. 
|   2|3   |   |   | 
|   7|   |   |   | 
|   |   |   |   | 
|____________|____________|____________|____________| 
|   |   |   |   | 
|   |   |   |   | 
|   |   |   |   | 
|____WG5_____|____WG6_____|____WG7_____|____WG8_____| 

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

Надеюсь, я понял вашу проблему. Для лучшего решения дополнительная информация о вашем алгоритме будет очень полезна.

Answer 2

Первый вопрос: «Требуется ли каждый рабочий элемент для синхронизации с другими?"

Поскольку, как правило, цифровые фильтры нужно только взаимодействовать с входными данными, чтобы обеспечить выход. Но они не рекурсивно использовать вывод предыдущих или соседей пунктов.

Если это так, вы не» t требуется какой-либо механизм синхронизации.

Во всех руководствах OpenCL должно быть БОЛЬШОЕ заявление об отказе от ответственности: «Локальная память работает быстро, но частная память еще быстрее, если вам не нужно использовать локальную память, потому что вы не нужно синхронизировать между потоками, а затем просто не использовать его ».

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

Если вы все еще хотите спуститься по этой дороге. Затем вам нужно сделать две части в вашем ядре: Fetch и filter. Каждая часть разделена барьером.

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

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

Answer 3

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

Из вашего описания проблемы (что, по общему признанию, я не полностью понимания), похоже, что любой данный рабочий элемент имеет зависимости от некоторого подмножества всех предыдущих рабочих элементов, в том числе и за пределами текущей рабочей группы. Если это так, у вас есть последовательная проблема, а не параллельная проблема. Если это так, вам нужно выяснить, как превратить его в параллельную проблему, удалив или, по крайней мере, уменьшив эти зависимости. Можете ли вы, например, ограничить количество «оглядываться назад» не выше некоторого числа? Если это так, вы можете вычислить параллельно элементы, которые далеко друг от друга, без зависимостей. Или, может быть, вы можете работать в параллельных различных наборах проблем (делайте только один рабочий элемент из каждого, но много параллельно).