Comments 65
(голос с галёрки fpga-дизайнеров)
интересно, как быстро это будет работать, если написать код на hdl-е достаточно прямыми руками…
интересно, как быстро это будет работать, если написать код на hdl-е достаточно прямыми руками…
Честно говоря, я не уверен, что это можно вообще написать на rtl за приемлемое время (скажем полгода). Говорю как RTL программист со стажем. Что до сравнения HLS и ручного кодирования, то уже давно вышла куча статей, где эти подходы сравниваются, и утрверждается, что HLS даже лучше делает. Потому что позволяет вам из одного C кода с помощью прагм получать разные архитектуры, в итоге вы можете выбрать в зависимости от ваших задач.
Один из примеров
Один из примеров
За полгода на RTL можно написать практически всё. Это я вам говорю как FPGA-дизайнер (не программист) со стажем.
Не знаю, что там за «куча статей», но за последние 15 лет я ни разу не видел, чтобы HLS давал лучшие (или сопоставимые) результаты, чем рационально написанный RTL, если речь не идёт про «синтез» путём подстановки готовых, вусмерть оптимизированных в RTL-е, библиотек.
Что HLS даёт результаты быстрее — видел. Что HLS даёт в несколько раз больший размер в кристалле — видел. Что HLS даёт частоту на 30-100% ниже — тоже видел. А вот чтобы HLS давал сопоставимые с RTL-ем результаты — увы. Но обратите внимание, я имею ввиду RTL, который как RTL, а не как С с подправленным под verilog синтаксисом.
Но я не буду спорить, что для быстрого получения результатов в ограниченном наборе областей HLS уже более-менее годен. Кстати, сколько приблизительно человеко-часов вы потратили на FPGA-часть этого проекта?
Не знаю, что там за «куча статей», но за последние 15 лет я ни разу не видел, чтобы HLS давал лучшие (или сопоставимые) результаты, чем рационально написанный RTL, если речь не идёт про «синтез» путём подстановки готовых, вусмерть оптимизированных в RTL-е, библиотек.
Что HLS даёт результаты быстрее — видел. Что HLS даёт в несколько раз больший размер в кристалле — видел. Что HLS даёт частоту на 30-100% ниже — тоже видел. А вот чтобы HLS давал сопоставимые с RTL-ем результаты — увы. Но обратите внимание, я имею ввиду RTL, который как RTL, а не как С с подправленным под verilog синтаксисом.
Но я не буду спорить, что для быстрого получения результатов в ограниченном наборе областей HLS уже более-менее годен. Кстати, сколько приблизительно человеко-часов вы потратили на FPGA-часть этого проекта?
Ок, я не готов обсуждать производительность RTL разработчиков, в любом случае есть гении, а есть посредственные.
Надо просто четко понимать область применимости HLS. Это дизайны с потоковыми вычислениями, длинными вычислительными конвейерами, дизайны где не критична latency, где регулярный доступ к памяти.
Основной код схемы я писал один примерно месяц-полтора в full time, это вместе с логическим верифицированием. Дальше были проблемы с таймингами, уже не связанные с HLS, а со структурой микросхемы, состоящей из 4 отдельных кристаллов, которые мы решали еще пару месяцев но уже в расслабленном режиме, переписываясь с инженерами Xilinx.
UPD. Да и приличный и доступный HLS транслятор Vivado HLS появился в 2013, а не 15 лет назад.
Надо просто четко понимать область применимости HLS. Это дизайны с потоковыми вычислениями, длинными вычислительными конвейерами, дизайны где не критична latency, где регулярный доступ к памяти.
Основной код схемы я писал один примерно месяц-полтора в full time, это вместе с логическим верифицированием. Дальше были проблемы с таймингами, уже не связанные с HLS, а со структурой микросхемы, состоящей из 4 отдельных кристаллов, которые мы решали еще пару месяцев но уже в расслабленном режиме, переписываясь с инженерами Xilinx.
UPD. Да и приличный и доступный HLS транслятор Vivado HLS появился в 2013, а не 15 лет назад.
Вопрос не в гениальности, а в подходе.
Если пишется RTL, то подразумевается, что вы сами выбираете архитектуру, которая реализует алгоритм, и имеете над ней полный контроль. И отличие RTL-дизайнера от программиста как раз в том, что он имеет опыт выбора правильной архитектуры, которая обеспечит нужные характеристики по скорости и размеру. Он может уметь это лучше или хуже, но он это умеет.
В случае же HLS-а вы этого лишены полностью или большей частью. И может наступить такой момент, когда вроде всё должно работать, но «не лезет» по тому или иному критерию. И очень хорошо, если есть возможность эту проблему обойти и весь проект не выкидывается в корзину со словами вида «скорость вычислений не удовлетворяет условиям заказчика».
Так что быстрое прототипирование при условии большого запаса по ресурсам — наверное, да. А что-то зажатое более строгими требованиями — пожалуй, нет.
Если пишется RTL, то подразумевается, что вы сами выбираете архитектуру, которая реализует алгоритм, и имеете над ней полный контроль. И отличие RTL-дизайнера от программиста как раз в том, что он имеет опыт выбора правильной архитектуры, которая обеспечит нужные характеристики по скорости и размеру. Он может уметь это лучше или хуже, но он это умеет.
В случае же HLS-а вы этого лишены полностью или большей частью. И может наступить такой момент, когда вроде всё должно работать, но «не лезет» по тому или иному критерию. И очень хорошо, если есть возможность эту проблему обойти и весь проект не выкидывается в корзину со словами вида «скорость вычислений не удовлетворяет условиям заказчика».
Так что быстрое прототипирование при условии большого запаса по ресурсам — наверное, да. А что-то зажатое более строгими требованиями — пожалуй, нет.
Не совсем так. Высокоуровневые средства все же весьма продвинулись за последнее время. Имеется большое количество настроек реализации. В RTL вы становитесь заложником выбранной «архитектуры» и если просчитались — вынуждены руками переписывать все с нуля. HLS же дает возможность перебирать те самые архитектурные решения, выбрав подходящее вам.
Слово «архитектура» вы заметили в моём комментарии. Но заметили ли вы слова «полный контроль»?
Выбрать правильную архитектуру, опираясь на достаточно конкретные расчёты и опыт — это и есть суть работы RTL-дизайнера. С тем же успехом можно назвать программиста заложником своего кода.
Не думайте, что я не знаком с возможностями HLS-а. Его возможности по выбору «архитектурных решений» весьма узки.
Вот, например, в обсуждаемой статье говорится что «Vivado HLS не могла создать схему, способную на одном и том же такте читать и писать один и тот же аппаратный массив». Когда я писал hdl-код, где была необходимость каждый такт читать-изменять-записывать значение некоего контекста хранящегося в BRAM-е по полностью произвольному адресу, я потратил на это от силы день. Применённое решение работало без потерь по скорости и было достаточно элементарным, но не таким, которое способен выбрать HLS.
Хорошо, когда есть обходные пути для решения различных ограничений HLSа. Но иногда этих путей нет или они приводят к весьма заметным потерям ресурсов или частоты.
Выбрать правильную архитектуру, опираясь на достаточно конкретные расчёты и опыт — это и есть суть работы RTL-дизайнера. С тем же успехом можно назвать программиста заложником своего кода.
Не думайте, что я не знаком с возможностями HLS-а. Его возможности по выбору «архитектурных решений» весьма узки.
Вот, например, в обсуждаемой статье говорится что «Vivado HLS не могла создать схему, способную на одном и том же такте читать и писать один и тот же аппаратный массив». Когда я писал hdl-код, где была необходимость каждый такт читать-изменять-записывать значение некоего контекста хранящегося в BRAM-е по полностью произвольному адресу, я потратил на это от силы день. Применённое решение работало без потерь по скорости и было достаточно элементарным, но не таким, которое способен выбрать HLS.
Хорошо, когда есть обходные пути для решения различных ограничений HLSа. Но иногда этих путей нет или они приводят к весьма заметным потерям ресурсов или частоты.
Да согласен, с записью и чтением на одном такте HLS показала себя не с лучше стороны. Еще у нас были большие проблемы с однотактовой реализацией «вихря Мерсенна», и только когда я нарисовал схему, я понял как написать С код так, чтобы он выдавал значения каждый такт.
С другой стороны мы использовали версию 2014.4, сейчас в 2016.x версии могло что-то улучшиться.
С другой стороны мы использовали версию 2014.4, сейчас в 2016.x версии могло что-то улучшиться.
Чуть быстрее, но на железе за существенно меньше денег. Что при трудозатратах на разработку может и дороже получиться.
В некоторые моменты между HLS-ом и RTL-ем как раз проходит граница помещается/не помещается в самый большой доступный кристалл. Или, например, в клоковый регион, как это произошло у автора статьи.
Поправка: не в клоковый регион, а в отдельный кристалл SLR. С пересечением клоковых регионов проблем не было.
Кстати а как это делается на практике? Там на все 4 кристалла объединенные клоковые сети от одного PLL (в пределах домена)? Или буферная синхронизация клоков?
Не совсем, не проблема использовать один и тот же клок в разных SLR, это никто не запрещает. Проблемы с цепями, которые пересекают границу SLR. На них очень большая задержка дополнительная получается. Их очень желательно пайплайнить с обеих сторон границы. На практике это делается так. Вы делаете разные модули и констрейните их размещение в разные SLR, и уже ручками в этих модулях ставляете регистры на интерфейсные сигналы. Так я и сделал, разделил код на два HLS ядра, законстрейнил в разные области и вставил регистры на уровне RTL. Но осталась проблема с основным ядром, которое ну никак в один SLR не влезало. Тут пришлось экспериментировать с разными прагмами во время синтеза в HLS.
Правильно ли я понимаю, что в итоге вы сравнивали производительность решения на CPU и решения на CPU + FPGA? Или же в последнем случае CPU не производил каких-либо значимых вычислений?
Тема сама по себе интересная, но что именно являлось целью работы?
Если именно исследование микротрубочек, то почему настолько маленькое количество частиц? Может я ошибаюсь, но судя по вики, микротрубочки подлиннее вообще должны быть. Понятно, что как кусочек посмотреть — можно и столько, но может больше — лучше будет?
Если показать использование FPGA для молекулярной динамики, то как я понял, перенести этот подход на другие задачи также эффективно слету не выйдет — используется фиксированое число соседей, ну и опять-таки обычно частиц на порядки больше.
Если именно исследование микротрубочек, то почему настолько маленькое количество частиц? Может я ошибаюсь, но судя по вики, микротрубочки подлиннее вообще должны быть. Понятно, что как кусочек посмотреть — можно и столько, но может больше — лучше будет?
Если показать использование FPGA для молекулярной динамики, то как я понял, перенести этот подход на другие задачи также эффективно слету не выйдет — используется фиксированое число соседей, ну и опять-таки обычно частиц на порядки больше.
Цели работы были две.
1. Получить реальное ускорение на реальной задаче. Маленькое число частиц — особенность задачи, потому что вычисляются изменения только на конце микротрубочки, а она еще растет или уменьшается (не учитывается в этой реализации, но вы можете посмотреть у нас на гитхабе, мы это допилили), так в любом случае нижняя часть стабильна и не интересна. Сейчас коллеги реально запускают вычисления и скоро опубликуют результаты.
2. Показать, что FPGA применимо в этой сфере, и не только в статьях зарубежных, но и у нас в России. Эта задача — первая ласточка. Скоро будут результаты по использованию метода Монте — Карло как для этой задачи, так и для других, например ряда финансовых задачь.
1. Получить реальное ускорение на реальной задаче. Маленькое число частиц — особенность задачи, потому что вычисляются изменения только на конце микротрубочки, а она еще растет или уменьшается (не учитывается в этой реализации, но вы можете посмотреть у нас на гитхабе, мы это допилили), так в любом случае нижняя часть стабильна и не интересна. Сейчас коллеги реально запускают вычисления и скоро опубликуют результаты.
2. Показать, что FPGA применимо в этой сфере, и не только в статьях зарубежных, но и у нас в России. Эта задача — первая ласточка. Скоро будут результаты по использованию метода Монте — Карло как для этой задачи, так и для других, например ряда финансовых задачь.
В работе показано, что реализация на FPGA быстрее CPU в 17 раз и быстрее GPU в 11 раз.
Тревожный звонок. Если топовая Tesla в таких задачах не опережает CPU как следует, то что-то пошло не так. Ну и выбор OpenCL для написания кода под Nvidia — это странно.
Мы честно говорим, что производительность GPU мала на этой задаче в первую очередь из-за того, что размер задачи мал, и в итоге загрузка GPU очень мала. Что поделать, такая задача. Мы и ищем в первую очередь такие, на которых можно конкурировать с GPU. OpenCL для GPU — это стандарт, который поддерживает Nvidia. Мы проверя на CUDA, было точно также. Дело в том, что OpenCL теперь работает и на FPGA.
А вы проводили профилировку кода на CPU/GPU? Есть ли аналогичные готовые решения (для CPU должны быть, по идее), которые показывают схожую с вашей производительность?
Иными словами: хотелось бы знать, что за 17-кратным превосходством стоит хорошее решение задачи на FPGA, а не меделенный код для CPU, который можно значительно ускорить. Я не занимался никогда разработкой под FPGA, но (по слухам :) ) это дело небыстрое — что будет, если потратить такое же время на написание оптимизированного кода с использованием intrinsic'ов и ассемблера?
Иными словами: хотелось бы знать, что за 17-кратным превосходством стоит хорошее решение задачи на FPGA, а не меделенный код для CPU, который можно значительно ускорить. Я не занимался никогда разработкой под FPGA, но (по слухам :) ) это дело небыстрое — что будет, если потратить такое же время на написание оптимизированного кода с использованием intrinsic'ов и ассемблера?
Роман, все что мы делали, я описал. Мы максимально старались сделать параллельный производительный код на CPU. Мы публикуем эти результаты открыто. Очень будем рады, если кто-то возьмется и попробует сделать быстрее как на CPU так и на GPU.
Но про ассемблер ничего не могу ответить, кажется что это сложно. Время разработки FPGA схемы я тут приводил в комментариях — 1.5 месяца отладка вычислительной схемы full time и два месяца расслабленной работы по укладываюнию во времянные ограничения.
Но про ассемблер ничего не могу ответить, кажется что это сложно. Время разработки FPGA схемы я тут приводил в комментариях — 1.5 месяца отладка вычислительной схемы full time и два месяца расслабленной работы по укладываюнию во времянные ограничения.
Это обычно показывает, что задача плохо ложится на классический SIMT парралелизм и содержит очень много dataflow обработки. Именно на этих сценариях, а также в ситуициях где много бренчинга, по идее FPGA и должен выигрывать.
«Загруженность ядер GPU cоставила 7% от одного мультипроцессора (SM)»
Это типа итого 0,47% общей мощности?
В общем, код не распаралеливали ни под ЦПУ, ни под ГПУ
Это типа итого 0,47% общей мощности?
В общем, код не распаралеливали ни под ЦПУ, ни под ГПУ
Расскажите, пожалуйста, побольше про модель, честно говоря, лениво читать [20-21]. Рассматривалась внеклеточная среда и голая физхимия, или клеточная, с учётом кэпирующих белков и прочей нечисти?
Насколько я, программист понял, модель такая. Есть структура молекул микротрубки, она шевелиться под действием сил взаимодействия между этими молекулами, а также под действием Броуноского движения. Это динамика. Каждый шаг динамики — это 0.2 нс. Ее мы и считали на FPGA. Еще есть кинетика, это вероятностьный процесс присоединения/отсоединения одного димера от трубочки. происходящий на временах порядка 1 мс (могу ошибиться ± один порядок). По сравнению с динамикой, кинетика — это процесс редкий, и поэтому считается на CPU.
Т.е. была среда, где шевелилась микротрубка, которая могла из окружающей среды присоединять свои молекулы, или те могли отрываться. Других белков вокруг не было. Вероятности присоединения/отсоединения задавлись параметрами, физическую суть которых я до конца тут не смогу изложить.
Т.е. была среда, где шевелилась микротрубка, которая могла из окружающей среды присоединять свои молекулы, или те могли отрываться. Других белков вокруг не было. Вероятности присоединения/отсоединения задавлись параметрами, физическую суть которых я до конца тут не смогу изложить.
удалось ли подобрать условия, в которых наблюдается:
— динамическая нестабильность (присоединение и уход мономеров на обоих концах)
— тредмиллинг (рост на одном конце с одновременным укорочением на другом)
— катастрофа с раскрытием МТ как цветка
Все три поведения описаны, но экспериментальное подтверждение условий, когда доминирует тот или иной тип роста, вроде пока не появилось. Уж очень трудоёмко и на грани разрешения аппаратуры.
Ещё такой вопрос: изучалось поведение одиночных МТ, или ещё дуплетов\триплетов?
— динамическая нестабильность (присоединение и уход мономеров на обоих концах)
— тредмиллинг (рост на одном конце с одновременным укорочением на другом)
— катастрофа с раскрытием МТ как цветка
Все три поведения описаны, но экспериментальное подтверждение условий, когда доминирует тот или иной тип роста, вроде пока не появилось. Уж очень трудоёмко и на грани разрешения аппаратуры.
Ещё такой вопрос: изучалось поведение одиночных МТ, или ещё дуплетов\триплетов?
— Моделировался только один конец
— с катастрофы с раскрытием мы начинали моделирование чтобы верифицировать наш генератор псевдослучайных чисел. Работали только с одной МТ. Вообще говоря, вопросы не ко мне, я могу вас в личке свести с учеными биофизиками, кто ставил задачу и вообще решают ее. Я только FPGA ускоритель =)
— с катастрофы с раскрытием мы начинали моделирование чтобы верифицировать наш генератор псевдослучайных чисел. Работали только с одной МТ. Вообще говоря, вопросы не ко мне, я могу вас в личке свести с учеными биофизиками, кто ставил задачу и вообще решают ее. Я только FPGA ускоритель =)
1. Поймать Фазли Иноятовича — это неочевидная задача, в точности по Гейзенбергу: или время, или место. Если и то, и другое — то надо бегать по 4 корпусам весь рабочий день. А других биофизиков понять лично мне значительно сложнее, большинство знает или что делали, или зачем. Мне проще Вас поспрошать, Вы-то уж точно с обеими группами общались.:) Ну и статьи я точно прочту, по этой и в сборнике ИСП РАН модель и конечная биологическая цель работы были непонятны.
2.
2.
с катастрофы с раскрытием мы начинали моделирование чтобы верифицировать наш генератор псевдослучайных чиселЭт как? Событие же очень редкое.
2. Ну не знаю, просто если задать начальные условия с уже сильно раскрытым венчиком, то трубочка раскрывается очень быстро. Посмотрите видео в этой статье, где моя коллега подробно описала, как мы реализоывали генератор псевдослучайный чисел.
Хотя может я и путаюсь, это не наверное не катастрофа все-таки, а просто динамическая релаксация…
Коллеги из группы Фазли сейчас как раз работают над продолжение моделирования на FPGA, я обязательно упомяну тут об их прогрессе.
Хотя может я и путаюсь, это не наверное не катастрофа все-таки, а просто динамическая релаксация…
Коллеги из группы Фазли сейчас как раз работают над продолжение моделирования на FPGA, я обязательно упомяну тут об их прогрессе.
Что-то ваша программа (из папки cpu) вылетает в VS2015 с ошибкой:
Run-Time Check Failure #2 — Stack around the variable 'long_d_t' was corrupted.
OpenMP в компиляторе не включал, если что.
Run-Time Check Failure #2 — Stack around the variable 'long_d_t' was corrupted.
OpenMP в компиляторе не включал, если что.
Ой, да у вас и вправду программа некорректная!
calc_grad_c принимает последним аргументом &long_d_t[i][j + 1], а на последней итерации этот ваш j+1 указывает за пределы массива!
Почините сначала код свой, а потом производительность измеряйте. Готов спорить, если прямыми руками с пониманием написать, то все ускорится в разы, особенно GPGPU версия.
calc_grad_c принимает последним аргументом &long_d_t[i][j + 1], а на последней итерации этот ваш j+1 указывает за пределы массива!
Почините сначала код свой, а потом производительность измеряйте. Готов спорить, если прямыми руками с пониманием написать, то все ускорится в разы, особенно GPGPU версия.
Насколько FPGA будет опережать GPU в классической молекулярной динамике? Стоит ли ждать отдельных устройств узкой направленности? Например прибор для аппаратного ускорения того же громакса. Насколько я понимаю под каждую отдельную задачу надо заново компилировать программу и зашивать в нее параметры потенциала межмолекулярного взаимодействия и прочие условия моделирования. Так ли это?
Эта задача решена институтом D.E. Shaw, они даже дальше пошли — кластер ASICов запилили (Anton, Anton II). Ускорение примерно на два порядка для первой версии ASIC.
Попробую составить более полную картину. Для молекулярной динамики у меня получается:
производительность одной FPGA=19.4 x 2 x Intel Xeon E5-2640 (2x6 cores), 92.2k atoms
производительность одной GeForce GTX 1080 = 11.4 x 2 x Intel Xeon E5-2650V3 (2x10 cores), 90.9k atoms
Кластер из 512 нод ASICов (Anton II) = 377 x GTX 1080, 23.6k atoms
Кластер из 512 нод ASICов (Anton II) = 941 x GTX 1080, 90.9k atoms
То есть производительность одной GTX 1080 примерно равна производительности одной FPGA. Кластер 512 ASICов производительней на 3 порядка. То есть один ASIC примерно равен по производительности одному GPU, но ASIC масштабируются «в лоб» намного лучше чем GPU. То есть можно взять 512 ASIC и получить производительность в 512 раз выше (в рамках одной симуляции), больше 512, вроде не получится уже. Производительность GPU масштабируется «в лоб» максимум раза в 2-4 ещё на данный момент.
Единственный смысл в Anton, на сколько я могу судить, — это неприрывные длинные симуляции порядка миллисекунд и более, которые на GPU не достижимы. Во многих случаях необходимости в этом нет, так как аналогичные задачи можно решить не «в лоб», а с помощью задействования кучи GPU и enhanced sampling алгоритмов типа REMD или Markov State Models, которые не требуют интенсивной коммуникации между GPU. Три порядка всё равно трудно догнать, но один-два порядка — вполне. Задействовать такие алгоритмы на Anton не выйдет — дорого по деньгам.
По цене кластер ASICов совершенно заоблачный (малосерийный). В открытую Anton не продают, но на конференциях ходили слухи что они готовы продать такой кластер за примерно $10 млн. Одна GTX1080 стоит 600$, в ноде будет порядка 700-800$ с учётом материнки, ЦПУ и т.п.
В плане гибкости, GPU, конечно, намного предпочтительней, как и в плане обслуживания. Софта для MD на GPU полно, в т.ч. бесплатного. Специалистов, знакомых с GPU, тоже полно (чего не скажешь про custom-made ASIC clsuter). Соответсвенно доступность разных «расширений» за пределами базовой стандарной МД на GPU намного лучше.
CPU/GPU
FPGA/CPU
ASIC cluster (Anton II)
производительность одной FPGA=19.4 x 2 x Intel Xeon E5-2640 (2x6 cores), 92.2k atoms
производительность одной GeForce GTX 1080 = 11.4 x 2 x Intel Xeon E5-2650V3 (2x10 cores), 90.9k atoms
Кластер из 512 нод ASICов (Anton II) = 377 x GTX 1080, 23.6k atoms
Кластер из 512 нод ASICов (Anton II) = 941 x GTX 1080, 90.9k atoms
То есть производительность одной GTX 1080 примерно равна производительности одной FPGA. Кластер 512 ASICов производительней на 3 порядка. То есть один ASIC примерно равен по производительности одному GPU, но ASIC масштабируются «в лоб» намного лучше чем GPU. То есть можно взять 512 ASIC и получить производительность в 512 раз выше (в рамках одной симуляции), больше 512, вроде не получится уже. Производительность GPU масштабируется «в лоб» максимум раза в 2-4 ещё на данный момент.
Единственный смысл в Anton, на сколько я могу судить, — это неприрывные длинные симуляции порядка миллисекунд и более, которые на GPU не достижимы. Во многих случаях необходимости в этом нет, так как аналогичные задачи можно решить не «в лоб», а с помощью задействования кучи GPU и enhanced sampling алгоритмов типа REMD или Markov State Models, которые не требуют интенсивной коммуникации между GPU. Три порядка всё равно трудно догнать, но один-два порядка — вполне. Задействовать такие алгоритмы на Anton не выйдет — дорого по деньгам.
По цене кластер ASICов совершенно заоблачный (малосерийный). В открытую Anton не продают, но на конференциях ходили слухи что они готовы продать такой кластер за примерно $10 млн. Одна GTX1080 стоит 600$, в ноде будет порядка 700-800$ с учётом материнки, ЦПУ и т.п.
В плане гибкости, GPU, конечно, намного предпочтительней, как и в плане обслуживания. Софта для MD на GPU полно, в т.ч. бесплатного. Специалистов, знакомых с GPU, тоже полно (чего не скажешь про custom-made ASIC clsuter). Соответсвенно доступность разных «расширений» за пределами базовой стандарной МД на GPU намного лучше.
CPU/GPU
FPGA/CPU
ASIC cluster (Anton II)
Спасибо за ответ. Мы исследуем мембранные процессы разделения в неравновесных условиях молекулярной динамикой и указанные алгоритмы нам не подходят. Мы как раз проводим длительное моделирование, один только выход на стационарность может занимать несколько месяцев реального времени на нескольких тысячах молекул.
Ну и как не прискорбно это признавать, в нашем университете компьютеры не считаются средствами научных исследований, и никаких ГТХ1080 на соответствующую статью расхода купить не получится ((. По факту приходилось часть ГПУ покупать за свои. Ладно хоть за немцев порадуюсь
Ну и как не прискорбно это признавать, в нашем университете компьютеры не считаются средствами научных исследований, и никаких ГТХ1080 на соответствующую статью расхода купить не получится ((. По факту приходилось часть ГПУ покупать за свои. Ладно хоть за немцев порадуюсь
Да за 10 миллионов можно отдельный суперкомпьютер собрать. А я уж было понадеялся, что FPGA это дешёвый и сердитый способ работать с MD.
дешёвый и сердитый способ работать с MD — это GPU. С несколькими GPU уже можно набрать материал достаточный для публикации в нормальных научных журналах.
D.E.Shaw research по сути спроэктировали и собрали special-purpose суперкомпьютер и специализорованный же процессор, отсюда и цена.
D.E. Shaw вообще заработал деньги (миллирды, похоже) на разработке специализированных суперкомпьютеров для финансов/биржи в 90е. Затем ушел из финансов и на заработанные деньги финансирует собственный исследовательский институт, занимается вычислительной биохимией в основном.
Ожидать что FPGA будет дешевле/быстрей GPU было излишне оптимистично — GPU в MD загружаются на полную, но при этом они массовые в производстве за счёт геймеров. Это подтверждается также тем, что кастомный ASIC разработанный специально для MD по производительности равен GPU.
D.E.Shaw research по сути спроэктировали и собрали special-purpose суперкомпьютер и специализорованный же процессор, отсюда и цена.
D.E. Shaw вообще заработал деньги (миллирды, похоже) на разработке специализированных суперкомпьютеров для финансов/биржи в 90е. Затем ушел из финансов и на заработанные деньги финансирует собственный исследовательский институт, занимается вычислительной биохимией в основном.Ожидать что FPGA будет дешевле/быстрей GPU было излишне оптимистично — GPU в MD загружаются на полную, но при этом они массовые в производстве за счёт геймеров. Это подтверждается также тем, что кастомный ASIC разработанный специально для MD по производительности равен GPU.
А если какой-нибудь NVIDIA Jetson TX1? Или кластер собрать из 128 Raspberry PI Zero?
На счёт Jetson точно не скажу, но raspberry не пойдет — слишком медленно. Что бы догнать один GPU, приходится 64 ядра Ксеонов задействовать, 128 малин такой производительности и близко не дадут, не говоря уж о совершенно недостаточной скорости интерконекта, в кластерах ЦПУ всякие супер производительные InfiniBand применяют, и то еле-еле хватает.
Я думал о GPU коде для своей задачи, но у меня не полноценная MD, так что я пока не уверен, что лучше использовать.
Создавать прибор для аппаратного ускорения MD на базе отдельного FPGA смысла, на мой взгляд, нет т.к. GPU не медленнее, удобнее, проще и уже всё готово и оптимизировано.
Безусловно нет смысла в попытках заменить GPU на FPGA путём эмуляции этого GPU на FPGA.
Но если есть возможность использовать гибкость FPGA, то всё может быть более интересным. Например, может быть интересной идея создать внутри FPGA объектное представление моделируемой системы, с нужными свойствами объектов и буквальными связями. Типа нейросети, хотя это может быть не самым удачным примером, т.к. нейросети из простых элементов уже научились моделировать на GPU весьма успешно.
Но если есть возможность использовать гибкость FPGA, то всё может быть более интересным. Например, может быть интересной идея создать внутри FPGA объектное представление моделируемой системы, с нужными свойствами объектов и буквальными связями. Типа нейросети, хотя это может быть не самым удачным примером, т.к. нейросети из простых элементов уже научились моделировать на GPU весьма успешно.
Главное преимущество FPGA — это возможность собрать длинный вычислительный конвейер, каждая ступень которого будет рабоать каждый такт над новыми данными, а сами данные последовательно проходить каждый такт через конвейер. Как было сделано в нашей реализации.
В случае МД всё упирается в количество FLOPS в конечном итоге, а оно зависит от транзисторного бюджета и частоты. По совокупности этих параметров GPU сильно опережает FPGA. Польза FPGA в том, что можно на полную загрузить все транзисторы там, где GPU полностью загрузить не получается. Но МД позволяет эффективно загружать GPU, так что я не вижу перспектив у FPGA, что подтверждается и соотв. попытками. Я уже давал ссылку, где люди закодили MD на 4 x Xilinx V5 и догнали один GPU. То же самое подтверждается тем, что даже команда довольно крутых спецов на многомиллионном бюджете смогла лишь догнать по производительности GPU, запилив специальный ASIC. Преимущество ASIC в том, что там можно более эффективный интерконнект организовать и улучшить масштабируемость, но не удельную производительность.
Тут что-то не так: на вашей GPU должно выполняться на порядок быстрее, чем на вашем CPU. Возможно, у вас расходятся потоки, или слишком много используется переменных и как результат они выделяются в глобальной памяти, или вы не используете память local, или вы недостаточно хорошо распараллелили свой алгоритм, или слишком часто используете барьер, или еще что-то.
Да, на всякий случай: OpenCL — хороший выбор для программирования GPU, он не медленнее CUDЫ.
Да, на всякий случай: OpenCL — хороший выбор для программирования GPU, он не медленнее CUDЫ.
Я думаю, что GPU плохо справляется из-за узкого горла к общей памяти, в которой находят массивы генератора псевдослучайный чисел.
Рандомные числа можно сгенерировать с помощью get_global_id() и побитовых операций. См пример тут: http://stackoverflow.com/questions/9912143/how-to-get-a-random-number-in-opencl
Я боюсь, что период таких генераторов будет недостаточным для задачи моделирования Броуновского движения. Мы реализовывали известный генератор "вихрь Мерсенна", имеющий период около 2^19937
В отдельной статье по генератору, на которую вы давали ссылку в тексте, в комментариях посоветовали весьма простой и эффективный PCG. Он вам по каким-то причинам не подошёл или просто не захотели менять собственноручно накоденное, пусть и в ущерб производительности?
Нас уже устраивало то, что было реализовано. У нас не было задачи перебрать множество генераторов и выбрать лучший для задачи, достаточно было найти первый работающий.
Чуть выше вы сами предположили, что ваш вариант генератора является узким местом в реализации на GPU из-за доступа к памяти. Т.е. вас устраивает, что GPU-реализация вероятно работает сильно медленнее чем это возможно?
Если честно, то в этом вашем ответе (и некоторых других ответах) проскальзывает нежелание потратить некоторое количество времени на оптимизацию вашего же (не обязательно лично вашего, но вашей команды) кода. При том что речь может идти не о процентах, а о разах ускорения. Выглядит немного странно.
Если честно, то в этом вашем ответе (и некоторых других ответах) проскальзывает нежелание потратить некоторое количество времени на оптимизацию вашего же (не обязательно лично вашего, но вашей команды) кода. При том что речь может идти не о процентах, а о разах ускорения. Выглядит немного странно.
Я вас не пойму. Генератор вирхь мерсенна изначально использовался в коде, который был у ученых. Мы попробовали сделать свой на основе центральной предельной теоремы — не вышло, был плохой период, трубочка разваливалась. В итоге запилили вихрь мерсенна, с ним все было ок. Эту же версию и портировали на GPU. Вы меня обвиняете в том, что я потратил мало времени на вылизование кода GPU? И даже не исследовал, как поведут себя другие генераторы на GPU.
Еще раз повторяю, у меня не было такой задачи.
Что плохого в том, что я хочу показать выгодную мне технологию в хорошем свете? Я же не говорю, что GPU всегда хуже, я говорю на данном алгоритме, который придумал не я, FPGA показывают себе лучше, вот и все.
Еще раз повторяю, у меня не было такой задачи.
Что плохого в том, что я хочу показать выгодную мне технологию в хорошем свете? Я же не говорю, что GPU всегда хуже, я говорю на данном алгоритме, который придумал не я, FPGA показывают себе лучше, вот и все.
Ок, возможно проблемы не у вас, а у того, кто ставил задачи перед вами и другими разработчиками.
Вы пишите:
Это некорректно. Быть может, было бы корректней сказать, что наивный порт на GPU в 11 раз медленнее вашей реализации на FPGA? Вы для демонстрации выбрали алгоритм генерации псевдослучайных чисел, который был разработан под CPU. Существуют алгоритмы (и реализации) для генерации псевдослучайных чисел для GPU, которые ничем не уступают Мерсенну и хорошо работают.
Это означает, что FPGA быстрее GPU в 11 раз.
Это некорректно. Быть может, было бы корректней сказать, что наивный порт на GPU в 11 раз медленнее вашей реализации на FPGA? Вы для демонстрации выбрали алгоритм генерации псевдослучайных чисел, который был разработан под CPU. Существуют алгоритмы (и реализации) для генерации псевдослучайных чисел для GPU, которые ничем не уступают Мерсенну и хорошо работают.
Я конечно не математик, но позвольте выразить свой скептис по поводу «был плохой период, трубочка разваливалась».
Скажите, а в чем вообще проблема подождать четыре секунды? Зачем вы мучаетесь с OpenMP, OpenCL, FPGA?
Наверное вам надо посчитать больше, чем одна трубочка? Раскройте, пожалуйста, а люди, возможно, подскажут как это лучше распараллелить.
Наверное вам надо посчитать больше, чем одна трубочка? Раскройте, пожалуйста, а люди, возможно, подскажут как это лучше распараллелить.
Sign up to leave a comment.
Применение FPGA для расчета деполимеризации микротрубочки методом броуновской динамики