Publish @Zeux: GDC Report – DirectCompute for Gaming: Supercharge your engine with Compute Shaders – Stephan Hodes and Layla Mah (AMD)

Задерживаюсь с постами про GDC, надо постить чаще. Буду догоняться!

People in this conversation:
Arseny Kapoulkine (Zeux) Add
CREAT Studios, Saber3D, Sperasoft (EA Sports – FIFA). Currently making kids happy at ROBLOX. http://zeuxcg.org/

Zeux: DirectCompute for Gaming : Supercharge your engine with Compute Shaders – Stephan Hodes and Layla Mah (AMD)
Поехали дальше! Доклад у которого был потенциал, но который абсолютно убит отсутствием внятного фокуса и стилем презентации…
Значит, первую половину докладывал мутный дядечка Stephan Hodes, там был такой DirectCompute tutorial, типа что писать в шейдере, что такое thread group, что такое shared memory, etc.
Мутный дядечка показывал сравнение в скорости чего-то без CS с чем-то с CS на графике с двумя столбиками, измеряющимися в FPS (типа 240 fps и 258 fps)
Я уж думал таких больше нет, а фиг там.

Кроме стандартного tutorial content была значит как use case предложена оптимизация separable фильтров на CS
Конкретно, утверждалось что при некоторых размерах фильтра можно выиграть в bandwidth
Методом “прочитаем данные которые нужны нескольким последовательно идущим пикселям в shared memory, а потом будем читать из нее”
Утверждается что есть бонус если bandwidth на shared memory в разы выше чем bandwidth в L1 texture cache.

Потом докладывалась Лейла, про DoF. Достаточно интересно, но ооочень быстро т.к. доклад был расчитан на целый таймслот по-моему, а досталась половинка.
Сначала вводная – рассказывала про то зачем нужен DoF и что главная проблема современных методов реализации DoF – блидинг из out of focus в in focus объекты
И вот значит есть такой алгоритм, diffusion DoF.
Про него уже несколько лет назад докладывались и в NVidia и в AMD, насколько я знаю там моделируется диффузия тепла на 2д изображении, где коэффициент диффузии зависит от circle of confusion (насколько точка в расфокусе)
И вот значит есть реализация от AMD на DX11 / PS, давайте ее как-нибудь ускорим на CS.

Дальше могут быть фактические ошибки, т.к. я рассказываю то что рассказали сегодня очень быстро, без знаний про сам алгоритм :)
Для того чтобы решить алгоритм, строится огромная СЛАУ, которая задается тридиагональной матрицей -
все коэффициенты нули, кроме двух или трех центральных диагоналей
Сторона матрицы изначально примерно равна размеру картинки, т.е. скажем миллион строчек.
Решив СЛАУ, задающуюся матрицей, мы получим выход DoF.
Решение СЛАУ состоит из двух частей – сначала мы за N итераций сводим ее до условно 1х1 матрицы
За каждую итерацию уменьшая кол-во неизвестных в два раза
Потом, чтобы посчитать результат, мы идем в обратную сторону, за каждую итерацию восстанавливая результат для промежуточной матрицы с какого-то шага из первого этапа.
В итоге для миллиона пикселей получается условно 20 проходов на уменьшение, и 20 проходов в обратную сторону.
Для экономии промежуточной памяти там первый шаг уменьшает не 2 а сразу в 4 раза, но это какая-то мелкая деталь…

Ну и вот значит есть реализация, которая делает 40+ draw calls с разными PS
Каждый следующий этап на входе принимает RT, заполненную предыдущим этапом
И генерирует RT на вход следующему этапу.
У AMD есть внутренняя тулза, про которую они думают как ее отдать в народ, которая собирает thread trace – активность всех compute тредов gpu (ну и наверное cpu там же).

Пособирав информацию этим тулом, стало очевидно, что реализация могла бы быть и побыстрее:
1. каждый шаг (которых 40+) это draw call, который читает из RT, записанной на предыдущем шаге – поэтому GPU ждет окончания записей предыдущего шага прежде чем запустить следующий
2. каждый шаг рисует full screen quad – пока он это делает, никакой больше полезной работы не происходит (т.к. окончания предыдущего DIP мы подождали)
3. количество тредов в wavefront на AMD 64, т.е. если мы обрабатываем меньше 64 элементов – мы тратим такты впустую. Часть из тех проходов обладали этим свойством.
4. каждый шаг пишет в RT, запись идет через ROP, ROP напрямую пишет в память – следующий шаг должен эти данные опять поднимать из памяти

Итак, предложенное решение, которое решает часть этих проблем
Пишется один CS, который сразу выполняет все 40+ шагов.
Поскольку шаги все равно зависимы друг от друга
То для синхронизации используем GroupMemoryBarrierWithThreadGroupSync или как-то так, короче __syncthreads() из CUDA
Синхронизация все равно есть, но она дешевле т.к. внутри CS а не сквозь весь GPU pipeline

Возникает несколько проблем.
Во-первых, все 80 RT (вместо RT кстати используются unordered access views) сразу привязать не получается т.к. есть лимит в D3D11
Поэтому все кладем в 1 огромный UAV

Во-вторых, совершенно непонятно, какой размер thread group ставить.
При выполнении PS пиксели группируются на векторы по 64 элемента сами
Включая запуск новых тредов по потребности в зависимости от количества пикселей итп
Поскольку разные этапы оперируют разными размерами данных
То было бы неплохо иметь разные thread group sizes для разных этапов
Но поскольку мы в пределах одного CS, то так не получается.
Решение – динамический параллелизм, а точнее эмуляция динамического distribution пикселей по тредам.
Пишется функция, которая по thread index и по размеру текущего обрабатываемого куска данных понимает, какой конкретно пиксель ей читать
Там есть всякие варианты, которые зависят от того, как устроен memory access, как сделать так чтобы доступ к памяти из разных тредов склеивался в один запрос к памяти итп
В ряде случаев оказывается что тред группа не нужна – т.е. элементов надо обрабатывать сильно меньше чем есть тред групп
Такие группы “отключаются” бранчингом
(если группа нужна то делаем работу иначе не делаем ничего)
Там были примеры кода с какой-то препроцессорной магией, чтобы все это работало, и результаты по скорости – все это было быстро промотано, вроде текущие результаты – порядка 50% прироста по сравнению с предыдущим кодом
И утверждается что результирующий CS код проще поддерживать

В качестве потенциальных pitfalls было названо GPR count – с шансами шейдер в котором все 40 стадий сразу скомпилируется в менее эффективный с точки зрения кол-ва используемых регистров шейдер чем большинство из оригинальных, что негативно скажется на способностях scheduling.
Наверное из вышеописанного мало что понятно, мне тоже было мало что понятно! Ждем слайдов, а у меня по этому докладу все.
Сложилось впечатление что тот мутный парень плохо понимает что такое GPU, а Лейла жжет.
Если бы она не навставляла на каждый слайд убогих powerpoint-овских анимаций, цены бы ей не было.

Исходный пейпр (который как я понял и улучшался) вот – http://amddevcentral.com/gpu_assets/An%20Optimized%20Diffusion%20Depth%20of%20Field%20Solver.ppsx (там еще в референсах предыдущая реализация, которую эта в свою очередь улучшает…)