Как оптимизировать код OpenCL для доступа к соседи?

Question

Редактировать: Результаты предлагаемых решений добавляются в конце вопроса.

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

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

Моя реализация OpenCL: я создаю буфер OpenCL 1 измерения, залейте его структурами, представляющими элементы, у которых есть «int neighbors [6]», где я храню индекс соседей в буфере. Я запускаю ядро, которое обращается к соседям и накапливает их информацию в переменные элемента, которые не рассматриваются на этом шаге, а затем я запускаю другое ядро, которое использует эти переменные для обновления элементов. Эти ядра используют только __global переменные.

Пример кода:

typedef struct{ 
    float4 var1; 
    float4 var2; 
    float4 nextStepVar1; 
    int neighbors[8]; 
    int var3; 
    int nextStepVar2; 
    bool var4; 
} Element; 

__kernel void step1(__global Element *elements, int nelements){ 
    int id = get_global_id(0); 
    if (id >= nelements){ 
    return; 
    } 
    Element elem = elements[id]; 

    for (int i=0; i < 6; ++i){ 
    if (elem.neighbors[i] != -1){ 
     //Gather information of the neighbor and accumulate it in elem.nextStepVars 
    } 
    } 
    elements[id] = elem; 
} 

__kernel void step2(__global Element *elements, int nelements){ 
    int id = get_global_id(0); 
    if (id >= nelements){ 
    return; 
    } 
    Element elem = elements[id]; 

    //update elem variables by using elem.nextStepVariables 
    //restart elem.nextStepVariables 
} 

Прямо сейчас, моя реализация OpenCL занимает в основном то же время, чем моя реализация C++.

Итак, вопрос: как бы вы (эксперты: P) решить эту проблему? Я прочитал о трехмерных изображениях, чтобы сохранить информацию и изменить шаблон доступа к окрестности, изменив NDRange на 3D. Кроме того, я прочитал о __local памяти, чтобы сначала загрузить все окрестности в рабочей группе, синхронизировать с барьером, а затем использовать их, чтобы уменьшить доступ к памяти.

Не могли бы вы дать мне несколько советов по оптимизации процесса, подобного тому, который я описал, и, если возможно, дать мне несколько фрагментов?

Thanx.

Редактировать: Третья и пятая оптимизации, предложенные Huseyin Tugrul, уже были в коде. Как уже упоминалось, here, чтобы заставить структуры правильно вести себя, они должны удовлетворять некоторым ограничениям, поэтому стоит понять, что для предотвращения головных болей.

Редактировать 1: Применение седьмой оптимизации, предлагаемой производителем Huseyin Tugrul, увеличилось с 7 до 60 кадров в секунду. В более общем эксперименте прирост производительности составлял около x8.

Редактировать 2: Применение первой оптимизации, предложенной Huseyin Tugrul, увеличилось примерно на x1.2. Я думаю, что реальный выигрыш выше, но скрывается из-за еще одного узкого места, которое еще не разрешено.

Редактировать 3: Применение 8-й и 9-оптимизаций, предложенные Huseyin Tugrul не изменилась производительность, из-за отсутствия значительного кода преимущества этих оптимизаций, стоит попробовать в других ядрах, хотя.

Edit 4: Передача инвариантных аргументов (например, n_elements или workgroupsize) к ядрам как #defines вместо арг ядра, как уже упоминалось here, повышение производительности около x1.33.Как поясняется в документе, это связано с агрессивной оптимизацией, которую компилятор может выполнять, зная переменные во время компиляции.

Edit 5: Применяя вторую оптимизацию, предложенную Huseyin Tugrul, но с использованием 1 бит на соседа и используя битовые операции, чтобы проверить, если сосед присутствует (так, если соседи & 1 = 0, верхний сосед присутствует, если! соседи & 2! = 0, бот-сосед присутствует, если соседи & 4! = 0, правый сосед присутствует и т. д.), увеличенная производительность в 1,1 раза. Я думаю, что это было главным образом из-за сокращения передачи данных, поскольку движение данных было и остается моим узким местом. Вскоре я попытаюсь избавиться от фиктивных переменных, используемых для добавления дополнений к моим структурам.

Edit 6: Устраняя структуры, которые я использовал, и создание разделенных буферов для каждого свойства, я устранил переменные отступы, экономя пространство, и был в состоянии optimice глобального доступа к памяти и локальное распределения памяти. Производительность увеличилась в 1,1 раза, что очень хорошо. Стоит делать это, несмотря на сложность и непроницаемость программ: P.

Answer 1

В соответствии с вашими шагами 1 и 2, вы не делаете работу с основным ядром gpu. Какова сложность вашего ядра? Каково ваше использование gpu? Вы проверили с программами мониторинга, такими как afterburner? Средние настольные игровые карты могут получать потоки 10 тыс., Каждый из которых выполняет итерации 10 тыс.

Поскольку вы работаете только с соседями, размер данных/размер вычислений слишком велик, и ваши ядра могут быть узким местом для vram bandiwdth. Ваш основной системный RAM может быть таким же быстрым, как и пропускная способность pci-e, и это может быть проблемой.

1) Использование выделенного кэша может быть причиной фактической ячейки ячейки потока в частных регистрах, что является самым быстрым. Затем соседи в __local массив, поэтому сравнения/вычисления выполняются только в чипе.

Ток нагрузки клеток в __private

нагрузки соседей в __local

начала зацикливания для локального массива

получить следующий сосед в __private от __local

вычислений

конца цикла

(если у него много соседей, строки после «Load neighbors in __local» могут быть в другом цикле, который поступает из основной памяти с помощью патчей)

Каков ваш gpu? Приятно, что это GTX660. У вас должен быть 64 кб управляемый кеш на единицу расчета. Процессоры имеют только регистры 1kB и не адресуются для операций с массивами.

2) Более короткая индексация может использовать один байт в качестве индекса соседнего хранилища вместо int.

Пример:

0=neighbour from left 
1=neighbour from right 
2=neighbour from up 
3=neighbour from down 
4=neighbour from front 
5=neighbour from back 
6=neighbour from upper left 
... 
... 

так что можно просто вывести индекс соседа из одного байта вместо 4-байтовой Int который уменьшает основную память доступа, по крайней мере, соседа доступ.Ваше ядро ​​будет выводить индекс соседа из верхней таблицы, используя его вычислительную мощность, а не мощность памяти, потому что вы сделаете это из основных регистров (__ privates). Если ваш общий размер сетки постоянный, это очень просто, например, просто добавление 1 фактического идентификатора ячейки, добавление 256 к id или добавление 256 * 256 к id или около того.

3) Оптимальный размер объекта может сделать размер вашей структуры/ячейки более кратным 4 байтам. Если ваш общий размер объекта составляет около 200 байт, вы можете поместить его или увеличить его с помощью пустых байтов, чтобы сделать ровно 200 байтов, 220 бит или 256 байтов.

4) Безветрный код может использовать меньше if-утверждений. Использование if-statement делает вычисления намного медленнее. Вместо того, чтобы проверять -1 как конец индекса нити, вы можете использовать другой способ. Бесполезный легкий сердечник не способен выдерживать тяжеловес. Вы можете использовать поверхностные-буферные ячейки для обертывания поверхности, поэтому вычисленные ячейки всегда будут иметь 6-соседей, поэтому вы избавитесь от if (elem.neighbors [i]! = -1). Стоит попробовать особенно для GPU.

Просто вычисление всех соседей выполняется быстрее, чем выполнение if-statement. Просто умножьте результат на ноль, если он не является допустимым соседом. Как мы можем знать, что это не действительный сосед? Используя байтовый массив из 6 элементов на ячейку (параллельно с массивом соседних идентификаторов) (invalid = 0, valid = 1 -> умножьте результат на это)

Оператор if находится внутри цикла, который рассчитывается для шесть раз. Развертка петли дает аналогичное ускорение, если рабочая нагрузка в контуре относительно проста.

5) Элементы данных Изменение порядка можно перемещать INT [8] элемент к самой верхней стороне структуры, чтобы доступ к памяти может стать более с получением таким образом меньшего размера элементов к нижней стороне можно прочитать в одной операции считывания.

6) Размер рабочей группы при попытке разного размера рабочей группы может дать производительность 2-3 раза. Начиная с 16 до 512 дает разные результаты.

7) Разделение переменных (только если вы не можете избавиться от if-statements) Разделение элементов сравнения из struct. Таким образом вам не нужно загружать целую структуру из основной памяти только для сравнения int или логического. Когда требуется сравнение, затем загружает структуру из основной памяти (если у вас уже есть локальная оптимизация памяти, тогда вы должны поместить эту операцию до нее, так что загрузка в локальный mem выполняется только для выбранных соседей)

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

8a) Magic Использование сдвига вместо деления по мощности 2. Выполнение подобных по модулю. Помещение «f» в конец плавающих литералов (1.0f вместо 1.0), чтобы избежать автоматического преобразования из double в float.

8b) Magic-2 -cl-mad-enable Опция компилятора может увеличить умножение + добавить скорость работы.

9) Скрытность с задержкой Оптимизация конфигурации выполнения. Вам нужно скрыть задержку доступа к памяти и позаботиться о занятости.

Get maximum cycles of latency for instructions and global memory access. 
Then divide memory latency by instruction latency. 
Now you have the ratio of: arithmetic instruction number per memory access to hide latency. 
If you have to use N instructions to hide mem latency and you have only M instructions in your code, then you will need N/M warps(wavefronts?) to hide latency because a thread in gpu can do arithmetics while other thread getting things from mem. 

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

11) Задержка Скрытие снова Попробуйте код ядра только с арифметикой (закомментируйте все мем доступ и начать их с 0 или sometihng вы хотите), то попробуйте код ядра с инструкциями доступа только памяти (закомментировать расчеты/сослагательного наклонения)

Сравнение времени ядра с исходным временем ядра. Что еще больше влияет на исходное время? Сконцентрируйтесь на том, что ..

12) Lane & Bank Конфликты Правильно ли конфликты LDS переулками и глобальные банковские памяти конфликты, поскольку одни и те же адреса accessings можно сделать в процессе замедления serialed пути (новые карты имеют транслируют способность уменьшить это)

13) Использование регистров Попробуйте заменить любых независимых местных жителей частными лицами, так как ваш GPU может использовать пропускную способность почти 10 ТБ/с с использованием регистров.

14) Использование регистров Не используйте слишком много регистров, иначе они будут разливаться в глобальную память и замедлить процесс.

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

16) Сбор Scatter Когда соседи разные частицы (например, в Nbody NNS) от случайных адресов памяти, его может быть трудно применить, но, build read Оптимизация может дать 2x-3x скорость перед оптимизацией (требуется оптимизация локальной памяти для работы), поэтому она считывает в порядке от памяти вместо случайного и переупорядочивает по мере необходимости в локальной памяти для разделения между (разбросом) и потоки.

17) Разделяй и властвуй Только в случае, когда буфер слишком велик, и копировать между хостом и устройством таким образом делает GPU ожидания простоя, а затем разделить его на две части, отправить их по отдельности, начинают вычисления, как только один приходит, отправить результаты возвращаются одновременно в конце. Даже параллелизм на уровне процесса может подтолкнуть gpu к своим пределам таким образом.

18)Пропускная способность из памяти квалификаторов. Когда для ядра требуется дополнительная пропускная способность «read», вы можете использовать ключевое слово «__constant» (вместо __global) для некоторых параметров, которые меньше по размеру и только для чтения. Если эти параметры слишком велики, вы можете по-прежнему иметь хороший поток из квалификатора '__read_only' (после квалификатора '__global'). Similary '__write_only' увеличивает пропускную способность, но в основном это производительность оборудования. Если это серия HD5000 от Amd, постоянная хорошая. Возможно, GTX660 быстрее работает с кешем, поэтому __read_only может стать более удобным (или Nvidia использует кеш для __constant?).

Имейте три части одного буфера с одним __global __read_only, один как __constant и один как __global (если их строить не штрафует больше, чем читает).

Только что протестировал мою карту с использованием примеров AMD APP SDK, полоса пропускания LDS показывает 2 ТБ/с, а постоянная - 5 ТБ/с (такая же индексация вместо линейной/случайной), а основная память - 120 ГБ/с.

19)Современные функции аппаратного трансцендентные быстрее, чем старый бит взломать (например, Quake-3 быстро обратными квадратный корень) версии

20)теперь есть OpenCL 2.0, который позволяет нерестилища ядра внутри так что вы можете дополнительно увеличить разрешение в 2-й точке сетки и разгрузить его в рабочую группу при необходимости (что-то вроде увеличения детализации завихрения на краях жидкости динамически)

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

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

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