В рамках рубрики «Задайте вопрос эксперту Intel» мы попросили ведущего специалиста Intel Константина Владимирова ответить на вопросы, связанные с гетерогенным программированием, тулкитом oneAPI и смежными интересными вещами. Результат превзошел все наши ожидания. Константин не пожалел времени и дал развернутые и обоснованные ответы, не опасаясь быть полемичным. По сути, получилась небольшая лекция о кросс-архитектурном программировании во всех его ипостасях: нюансах оффлоада, оптимизациях, стандартах и так далее.
Передаем микрофон эксперту. Ну а комментарии отданы аудитории.
Вопрос Soarex16 Насколько трудоёмким будет переход с OpenCL на oneAPI и какие преимущества от этого можно будет получить?
Ответ. Переход на DPC++ может быть сложным, но, с моей точки зрения, он того стоит. Здесь есть два основных этапа.
Во-первых, это переход с вашего языка гетерогенного программирования (OpenCL, Vulkan compute), который, скорее всего, основан на API. Тут у вас фора в том, что вы уже знаете предметную область, а сложность в том, чтобы переключить мышление с непосредственного управления через API на чуть более неявные языковые конструкции.
Во-вторых, это переход с вашего языка написания хостовых программ. Если вы всю жизнь оффлоадили с чистого C, то входной порог равен порогу перехода с C на C++, а это довольно высоко.
Ради чего же стараться?
Во-первых, DPC++ отлично выполняет за программиста рутинную работу. Вы очень быстро забудете как страшный сон все эти явные вызовы clXXXYYY, и что там значит шестой аргумент, и не забыли ли вы код возврата. Многие объектно-ориентированные врапперы прячут рутину не хуже, но обычно ценой перехода со стандартного API OpenCL на не такое уж и стандартное API враппера (я видел те ещё велосипеды). В случае DPC++ вы просто пишете стандартный SYCL с расширениями Intel (которые, возможно, тоже скоро станут стандартным SYCL).
Во-вторых, DPC++ предусматривает совместную компиляцию, то есть вы можете быть уверены в типах и у вас не будет проблем по границами API с размерами, паддингом, выравниванием. Вы пишете ядро и хостовый код в одном файле, и это один и тот же код. С использованием USM вы также сможете намного проще работать со сложными структурами данных.
В-третьих, DPC++ это настоящий C++, то есть он допускает обобщённое программирование. Например, простейшее ядро для сложения двух векторов:
auto kern = [A, B, C](cl::sycl::id<1> wiID) {
C[wiID] = A[wiID] + B[wiID]; // какого типа A, B и C? Какого угодно!
};То же самое на OpenCL:
_kernel void vector_add(__global int *A, __global int *B, __global int *C) {
int i = get_global_id(0);
C[i] = A[i] + B[i]; }Вы видите, что я вынужден был указать на OpenCL тип int. Если мне нужен будет float, я должен буду либо писать другое ядро, либо использовать препроцессор, либо внешнюю кодогенерацию. Получить в своё распоряжение почти все возможности C++ может быть немножко страшно, если у вас не было опыта с C++. Но это обычное дело, когда речь идёт о серьёзном технологическом сдвиге.
И все преимущества этим не исчерпываются. Я кое-что ещё упомяну в следующих ответах.
Так что я бы на вашем месте уже скачал компилятор и попробовал, благо с пакетом OneAPI это сделать совсем несложно.
Вопрос JusterБудут ли каким-то образом связаны OpenVINO и oneAPI?
Ответ. Сейчас дистрибутив OpenVINO — это часть поставки OneAPI. Обучение нейросетей и их использование — это вычислительно тяжёлые задачи, которые сильно выигрывают от гетерогенного программирования. Я полагаю, рано или поздно все компоненты OneAPI будут давать возможность использовать все доступные вам вычислительные ресурсы: и графические ускорители, и специальные ускорители вроде Nervana и FPGA. И всё это не выходя из языковой парадигмы и системы типов вашей программы на C++.
Вопросы из почтыСтараюсь понять, как будет выглядеть AI hardware accelerator через 3 года, помогите пожалуйста в этом. Есть интересная компания Graphcore и её IPU — это устройство не менее эффективно, чем FPGA, но его намного проще программировать — Python с поддержкой TensorFlow и других фреймворков. Получается, что, если обещания Graphcore будут выполнены, необходимости в FPGA на рынке машинного обучения не будет т.к. для datascientists Python намного удобнее, чем C++.
Вы согласны с тем, что FPGA в сравнении с решениями программируемыми на Python не подходит рынку машинного обучения? Если рынок ML потерян, какие еще широкие применения FPGA вы видите?
В каких применениях вы видите неизбежную необходимость гетерогенного программирования, где нельзя обойтись более удобными инструментами вроде Python?
Ответ. Я бегло поглядел, что это за IPU. Ещё одна железка, на которую все будут оффлоадить. Эти ребята конкурируют с GPU и со специальными ускорителями, а не с FPGA.
В задачах, под которые специализированная железка заточена, она всегда побьёт FPGA, например, рендерить видео лучше на видеокарте и т.д. Но в мире (в том числе в мире ML) есть масса задач, для которых ничего специального не придумано или не выпущено, и тут FPGA останутся незаменимыми всегда. Например, потому что есть вопрос цены и, чтобы быть дешёвой, специализированная железка должна быть массовой.
Допустим теперь, что указанный IPU действительно крут. Это не отменит гетерогенное программирование, наоборот, наличие такого отличного ускорителя подхлестнёт его. И в том числе даст гигантскую фору OneAPI и DPC++, потому что рано или поздно кто-нибудь скажет «я хочу пользоваться и вашим IPU, и моим GPU из одной программы». Скорее рано, потому что гетерогенное программирование именно об этом. Его смысл — это оффлоад подходящей задачи на подходящее устройство. Задача может исходить откуда угодно. И это устройство может быть чем угодно, оно может быть даже тем же самым устройством, на котором выполняется программа. Например если вы оффлоадите ядро, написанное на ISPC и максимально утилизирующее векторные возможности Xeon, вы можете оффлоадить его сами на себя и всё равно быть в значительном выигрыше. Главный критерий тут — это производительность. Ну а производительности в этом мире никогда не будет слишком много. Даже с самыми лучшими в мире ускорителями.
Что до Python и его удобства… Признаюсь сразу, я не люблю динамически типизированные языки: они медленные и вместо нормальной ошибки компиляции там надо ждать два часа до падения в рантайм из-за неверного типа. Но я не вижу, чем плохо делать те же оффлоады из-под Python. К слову, в OneAPI уже сейчас входит крайне удобный по разным отзывам Intel Distribution for Python.
То есть в мире мечты любителей Python вы пишете на нем программу и оффлоадите с неё на все ускорители, которые можете найти, используя OneAPI, а не кучу вендор-специфичных библиотек. Другое дело, что при таком подходе вы упускаете сквозную типизацию и возвращаетесь в крайне ненадёжный мир API-based программирования. Возможно, развитие DPC++ подтолкнёт коммьюнити к более активному использованию более подходящих средств, таких как C++.
Вопрос из почтыПроизводительность в сравнении с OpenCL. Должны же быть налоги на роскошь – т.е. накладные расходы. Есть ли замеры?
Ответ. В интернете можно найти массу замеров с самыми разными результатами, в зависимости от компилятора, задачи и качества реализации. Я в качестве личного исследования мерил на простых задачах (SGEMM, DGEMM) на моём ноутбуке (интегрированная графика Skylake), и я видел, что пока что некоторая просадка есть (в пределах процентов). Но мне кажется, это следствие того, что всё это пока что бета.
В теории должно быть в итоге ускорение, а не замедление, то есть в принципе вся эта роскошь должна иметь отрицательную стоимость. Всё дело тут в компиляторе. Когда у вас программа состоит из единого исходника и обрабатывается как единая программа, компилятор получает фантастические, невероятные возможности к оптимизации: вынесение общего кода, разворот циклов, перестановка участков кода и всё остальное, что компилятор просто не может сделать в API-based подходе, но рано или поздно обязательно научится при модели единого исходника.
К тому же DPC++ будет иметь отрицательную стоимость и с точки зрения времени разработки. Простой пример — это SYCL accessors, которые компилятор уже сейчас использует для расстановки событий и управления асинхронными очередями.
deviceQueue.submit([&](cl::sycl::handler &cgh) {
auto A = bufferA.template get_access<sycl_read>(cgh);
auto B = bufferB.template get_access<sycl_read>(cgh);
auto C = bufferC.template get_access<sycl_write>(cgh);
....
deviceQueue.submit([&](cl::sycl::handler &cgh) {
auto A = bufferA.template get_access<sycl_read>(cgh);
auto B = bufferB.template get_access<sycl_read>(cgh);
auto D = bufferD.template get_access<sycl_write>(cgh);
Тут компилятор видит, что обе посылки только читают A и B и пишут независимые буферы C и D, в итоге он видит возможность отправить их параллельно, если глобальных размеров хватает.
Конечно, педантично написанная программа на OpenCL может делать это так же хорошо, но затраты времени на разработку при нетривиальном ядре будут несопоставимы.
Вопрос из почтыАктуальны ли все способы оптимизации OpenCL приложений для DPC++? Что нового к ним добавится?
Ответ. Я бы сказал, что большую часть тонкой ручной оптимизации, которая сейчас делается писателями ядер, может и должен делать компилятор. Точно так же я, например, считаю вредной практикой ручную простановку инлайн-ассемблера в программах на C++, потому что, даже если она даёт тактический выигрыш, она мешает оптимизациям и при развитии и переносе продукта выступает как негативный фактор. Ну вот и OpenCL — это теперь тоже ассемблер.
Что до более подробного ответа, то я боюсь тут бездны. Вот, например, есть известный документ Intel «OpenCL Developer Guide for Intel Processor Graphics». И там есть раздел о том, как надо стараться, чтобы не поставить где лишнюю синхронизацию.
Так вот, с моей точки зрения, это нечеловеческая задача в принципе. Люди крайне плохи в рассуждениях о многопоточной синхронизации и склонны лепить синхронизацию либо консервативно, либо неправильно, либо то и другое сразу — я примерно так расставляю запятые (но мы это исправили — прим. редакции).
С другой стороны, в DPC++ вместо написания кода с явными барьерами вроде вот такого:
for (t = 0; t < numTiles; t++) {
const int tiledRow = TS * t + row;
const int tiledCol = TS * t + col;
Asub[col][row] = A[globalRow * AY + tiledCol];
Bsub[col][row] = B[tiledRow * BY + globalCol];
// Synchronise to make sure the tile is loaded
barrier(CLK_LOCAL_MEM_FENCE);
// .... etc ....Вы, скорее всего, напишете явную итерацию parallel_for_work_group, внутри которой group.parallel_for_work_item
cgh.parallel_for_work_group<class mxm_kernel>(
cl::sycl::range<2>{BIG_AX / TS, BIG_BY / TS},
cl::sycl::range<2>{TS, TS}, [=](cl::sycl::group<2> group) {
// .... etc ....
for (int t = 0; t < numTiles; t++) {
group.parallel_for_work_item([&](cl::sycl::h_item<2> it) {
// .... etc ....
Asub[col][row] = A[globalRow][tiledCol];
Bsub[col][row] = B[tiledRow][globalCol];
});
// тут компилятор поставит синхронизацию локальной памяти, если она нужнаВ итоге вам вообще не надо будет ставить синхронизацию руками, и весь раздел можно выкидывать.
И так можно пройтись во всем разделам. Что-то выживет, что-то уйдёт. Я предвижу появление нового документа «Оптимизация под DPC++», но должно пройти время, так как все реально работающие техники нарабатываются только потом и кровью
Вопрос из почтыВ OpenCL существует ограничение – в ядре нельзя использовать «далекие данные», т.е., например, реализовать “широкий фильтр”, использующий в одном вычислении входные данные с большой группы пикселей, превышающей по размеру рабочую группу OpenCL. Что в этом отношении предлагает DPC++?
Ответ. Ну, так уж и нельзя. Я конечно особо не пишу кернела… Но совершенно точно там можно использовать всю глобальную память какая есть, просто надо следить за тем, чтобы работать атомарными операциями (или внешне синхронизировать иерархические кернела). А можно ещё и System SVM подключить (ну, или там USM в DPC++).
Увы, всё это крайне неэффективно, и я не люблю все эти трюки. К тому же их сложно оптимизировать компилятором.
И вот, если говорить о прямых и эффективных решениях, то, разумеется, в DPC++ нет никакой магии. Ваша программа в итоге всё равно разделяется на части: хостовый код и код девайсов, и на код девайсов влияют все ограничения девайсов. Максимальные размеры рабочей группы — это тот настоящий параллелизм, на который способно ваше железо. Всё что поверх этого — это просто способы выкрутиться, драматически негативно влияющие на производительность. Именно поэтому DPC++ предоставляет возможность сделать вот так: device.get_info<sycl::info::device::max_work_group_size>() и дальше уже решать как жить с полученным числом.
Соблазнительно, конечно, было бы сделать в DPC++ модель, когда программист работает как угодно с циклами любой длины, а компилятор смотрит, что с этим дальше сделать, но это было бы смертельно неправильно, потому что прятало бы константы, а то и асимптотику добавленной сложности вычислений, появляющейся ниоткуда. Александреску по другому поводу писал, что «инкапсуляция сложности должна считаться преступлением», и тут это тоже применимо.
Иногда помогает пересмотр самого алгоритма. Тут DPC++ делает вещи проще, потому что более структурный код проще рефакторить. Но это так себе утешение.
Вопрос из почтыDPC++ основан на SYCL. Но что, если залезть глубже под капот, какие отличия от OpenCL в реализации back end, если они вообще есть. Например, механизм распределения работы между гетерогенными устройствами – тот же, что и у OpenCL?
Ответ. Если залезть под капот, то это и есть OpenCL. Все преимущества и сильные стороны SYCL — это преимущества и сильные стороны языка, то есть фронтенда. Из-под фронтенда идёт старый добрый SPIRV который заходит на бэкенд и там оптимизируется (часто уже во время исполнения, т.е. JIT'тится) уже под конкретную видеокарточку точно так же, как под неё бы оптимизировался OpenCL.
Другое дело, что механизм распределения работы между гетерогенными устройствами — это как раз больше фронтенд, чем бэкенд, потому что именно хостовый код решает, что куда и как направить. А хостовый код получается из DPC++. Я уже чуть выше показал пример, как компилятор может сам на основании аксессоров принять решение о распараллеливании посылок. И это только верхушка айсберга.
Вопрос из почтыБиблиотеки. Да, мы не говорим про CUDA. Но знаем, что для разработчиков CUDA есть очень полезные библиотеки, работающие с высокой производительностью на GPU. OneAPI тоже кое-какие библиотеки содержит, но, например, IPP – архиполезной штуки для работы с изображениями, в oneAPI/OpenCL нет. Будет ли что-то, и как в этом случае переходить с CUDA к oneAPI?
Ответ. Переход с CUDA на единый открытый стандарт будет сложным, но неизбежным. Разумеется, сейчас у CUDA более зрелая инфраструктура. Но особенности её лицензирования — это блокирующий недостаток, потому что на рынке гетерогенных систем появляется всё больше игроков, всё больше интересных карт и ускорителей от разных производителей.
Разношёрстность существующих API делает использование этого мира возможностей сложным для программистов с опытом классического CPU. Что и приводит к OneAPI или чему-то вроде него. Тут магия не в рывке Intel в графику, а в том, что Intel открывает двери в DPC++ всем желающим. Мы даже не владеем стандартом SYCL, он принадлежит Khronos group и все расширения Intel — это расширения в Khronos, куда может коммититься кто угодно (и там есть представители ото всех крупных игроков). А значит появятся (уже появляются) и библиотеки и коммьюнити, и куча вакансий в этом направлении.
И конечно IPP перепишут под новые реалии. Я не имею отношения к IPP, но использование DPC++ — это здравый смысл, а там сидят здравомыслящие люди.
Но важнее то, что сейчас как раз тот момент в истории, когда вы можете написать свою библиотеку, которая превзойдёт IPP и которой будет потом пользоваться весь мир. Потому что открытые стандарты всегда побеждают.
Вопрос из почтыЕсли сравнить запуск алгоритмов нейронных сетей Training и Inference на Nervana и FPGA — какие будут отличия в программировании и полученной эффективности?
Ответ. Ничего не знаю про детали программирования FPGA, я компиляторы пишу. Но у меня встречный вопрос. А как будем сравнивать? На стандартных бенчмарках неспортивно, Nervana под них вылизывалась. Но в случае, если у вас нечто интересное, то FPGA развязывает вам руки, а положить это нечто на Nervana может быть долго, дорого, всё вот это.
Получается, что сам вопрос как бы из серии «кто сильнее слон или кит». Но это не настоящий вопрос. Настоящий вопрос звучит так: как впрячь и слона и кита в одну тележку? Ну или хотя бы распределить, скажем, чтобы слон тянул её по суше, а кит по морю.
В случае OneAPI у вас будет одна и та же программа на, в общем, стандартном C++. И вы можете её сами написать и запустить с оффлоадом туда и туда. Это будет та самая задача, которая вас интересует, на которой вы сами сможете померить и оптимизировать производительность. Единый стандарт и единый интерфейс к гетерогенным устройствам как раз и будет шагом к тому, чтобы в таких вопросах сравнивать яблоки с яблоками.
Например так: «что лучше для %моей задачи% с точки зрения простоты программирования и эффективности — вынести вот эту часть на FPGA, а вот эту оставить на Nervana или вот эту часть разбить на две, а вот эту переписать для GPU?»
И вся история с OneAPI — она как раз для того, чтобы вы сказали «а чего тут долго думать, я сейчас быстренько попробую, ЭТО ЖЕ ПРОСТО».
Пока что нет, не просто. Но будет.
Послесловие от эксперта
Всем спасибо за вопросы. Возможно и даже вероятно, что я был неправ, неточен и делал ошибки. Это бывает, в интернете постоянно кто-то неправ.
Надеюсь, что я смог кого-нибудь заинтересовать гетерогенным программированием и DPC++. Хочу всем порекомендовать сайт sycl.tech, где лежат тонны докладов, в том числе и от экспертов с мировым именем (требуется английский)
Всем добра!
P.S. от редакции. На этот раз единогласным решением редакции приз за лучший вопрос было решено вручить… автору ответов. Думаю, вы согласитесь, что это справедливо.
