Перед вами введение в программирование на языке Data Parallel C++ или, коротко, DPC++. DPC++ основан на Khronos SYCL — это означает, что перед нами модель современного параллельного программирования. Новейшим текущим стандартом Khronos является SYCL 1.2.1, хотя предварительная спецификация SYCL 2020 уже доступна для изучения. Intel и другие участники рабочей группы SYCL в настоящее время занимаются финализацией следующей версии спецификации. DPC++ содержит расширения, которые облегчают использование SYCL, при этом многие из них, как ожидается, войдут в состав SYCL 2020. Внедрение таких расширений в компилятор DPC++ помогает сообществу оценить их эффективность заранее перед стандартизацией.
Для кого предназначена эта статья
Эта статья — для программистов, уже имеющих достаточное представление о С++ и параллельности. Изучение С++ и параллельности — непростая задача, и написано об этом немало. О SYCL информации намного меньше, а о DPC++ — еще меньше, поэтому на них мы и сосредоточимся.
SYCL имеет свои корни в OpenCL, их модели исполнения довольно схожи. Если вам нужна помощь в понимании модели исполнения SYCL/OpenCL, обратитесь к этому обзору-презентации(на английском, но кратко, ёмко и с картинками).
Для кого НЕ предназначена эта статья
Когда я рассказываю о SYCL, то часто говорю: «Если вам нравится современный С++, то понравится и SYCL, потому что это определенно современный С++». И наоборот, если вам не по душе С++, то не понравятся и SYCL с DPC++. Поэтому, если вы не хотите писать на современном С++, эта учебная статья не для вас.
OpenMP 5.0 имеет почти тот же функционал, что и SYCL/DPC++, но поддерживает «большую тройку» ISO-стандартизированных языков: С, С++ и Fortran. Если вы хотите программировать на CPU и GPU с использованием Fortran, C или C++ до 11 версии с применением открытого промышленного стандарта, выберите OpenMP.
Другой альтернативой SYCL/DPC++ без C++ является OpenCL. OpenCL намного более многословен, чем SYCL, но если вы программист на С, то, скорее всего, предпочтете явный контроль за эффективностью синтаксиса.
Сложение векторов в SYCL
Мы начнем со сложения векторов, которое в некотором смысле является аналогом «Hello, world!» в мире HPC и численных методов. Очевидно, вывод символьной строки не может считаться хорошей задачей для языка параллельного программирования.
Операцией, которую мы пытаемся реализовать, является SAXPY — A умножить на X плюс Y, одинарной точности. На C (или C++ в простейшем случае) это делается так:
Хотя, надо отметить, что существует много способов написать этот код на С++. Например, мы можем использовать диапазоны, при этом код будет выглядеть более похожим на грядущий SYCL вариант. Но в данной статье рассказ о всех возможных вариантах написания цикла на С++ не предусмотрен, сосредоточимся на SYCL.
Вот как выглядит тот же цикл в SYCL. Здесь надо много объяснять, поэтому разобьем код на маленькие кусочки.
Как вы, наверное, догадались,
parallel_for — это параллельный цикл for. Тело цикла представляет собой лямбда-выражение. Лямбда выражения — это безымянные локальные функции, которые можно создавать прямо внутри какого-либо выражения. Их код выглядит как [..]{..}.Оператор цикла выражен в терминах
sycl::range и sycl::id. В нашем простом примере оба они одномерны, что указывается с помощью <1>. В SYCL диапазоны и, соответственно, идентификаторы могут быть одно-, двух-, или трехмерными. OpenCL и CUDA имеют то же ограничение.Возможно, для вас непривычно записывать циклы таким образом, но это согласуется с тем, как работают «лямбды». И, конечно, если вы когда-либо использовали параллельный STL, TBB, Kokkos или RAJA, то узнаете этот шаблон.
Возможно, вас удивит аргумент
<class saxpy> в цикле parallel_for. Это просто способ дать наименование ядру, что необходимо, поскольку вы можете захотеть использовать разные С++ компиляторы SYCL на хосте и устройстве. В этом случае два компилятора должны договориться об имени ядра. Во многих компиляторах SYCL, таких как Intel DPC++, этого не требуется. И мы можем попросить компилятор не беспокоиться о поиске имен с помощью опции -fsycl-unnamed-lambda.Сейчас мы не будем останавливаться на том, что значит
h в h.parallel_for. Вернемся к этому позже.Очереди в SYCL
Одна из главных проблем в гетерогенном программировании — различные типы элементов, обрабатывающих данные, и, зачастую, различные типы памяти. Это делает компиляторы и рантайм более сложными. Программная модель SYCL, конечно же, содержит в себе гетерогенное исполнение, хотя и реализованное на более высоком уровне, чем в OpenCL, а иногда — выраженное неявно. В отличие от других популярных программных моделей GPU, ядра SYCL могут быть встроены в код хост-программы, что повышает читаемость и удобство создания.
Если мы хотим сделать вычисления на устройстве, нужно создать рабочую очередь:
По умолчанию селектор указывает на GPU при его наличии, и на CPU в противном случае. Мы можем создать очереди, ассоциированные с конкретным типом устройства, таким образом:
Выбор селектора
host_selector() или cpu_selector() может существенно влиять на результат, даже если они показывают на одно и то же железо, поскольку host_selector() может использовать последовательное исполнение для целей отладки, при этом cpu_selector() использует рантайм OpenCL и задействует все имеющиеся процессорные ядра. Кроме того, JIT компилятор OpenCL в этом случае может генерировать другой код, поскольку он использует совершенно другой компилятор. Не стоит думать, что из-за того, что хостом всегда является CPU, хост и CPU означают одно и то же в SYCL.Управление данными в SYCL с использованием буферов
Классическим способом управления данными в SYCL являются буферы. Буфер SYCL — это непрозрачный контейнер. Выглядит красиво, но некоторые приложения хотели бы использовать указатели, предоставляемые USM расширением, — об этом позже.
В данном примере пользователь размещает С++ контейнер на хосте и затем передает его в SYCL. Пока не применен деструктор SYCL буфера, пользователь не может получить доступ к данным через не-SYCL механизм. Аксессоры SYCL, о которых пойдет речь далее — важный аспект управления данными буферов.
Контроль исполнения устройства
Поскольку код устройства может потребовать другой компилятор или механизм генерации кода, важно ясно обозначить секции кода устройства. Ниже мы видим, как это выглядит в SYCL 1.2.1. Мы используем метод
submit чтобы добавить задачу в очередь устройства q. Данный метод возвращает opaque handler, согласно которому мы исполняем ядра, в данном случае посредством parallel_for.
Мы можем синхронизировать исполнение устройства с помощью метода
wait(). Существуют более тонкие способы синхронизировать исполнение устройства, но мы начнем с самого простого и грубого.Некоторые пользователи могут посчитать код выше несколько громоздким, особенно в сравнении с моделями типа Kokkos. Компилятор Intel DPC++ поддерживает сжатый синтаксис.
Вычислительные ядра и буферы
Аксессоры — это финальный фрагмент нашей первой SYCL программы. Аксессоры, возможно, незнакомы программистам под GPU, но у них есть несколько преимуществ по сравнению с другими методами. В то время как SYCL позволяет программисту перемещать данные принудительно, с использованием, например, метода
copy(), методы аксессора не требуют этого, поскольку они генерируют граф потоков данных, который может использоваться компилятором или рантаймом для передвижения данных в нужное время. Это особенно эффективно, когда множество ядер задействуются последовательно. В данном случае SYCL-реализация заключит, что данные используются повторно и не вернет их без нужды на хост. Также мы можем запланировать асинхронное перемещение данных (т.е. накладывающееся на исполнение устройства). Опытные GPU-программисты могут сделать это вручную, но, зачастую аксессоры SYCL показывают лучшую производительность по сравнению с программами на OpenCL, где данные приходится перемещать принудительно.
Однако, поскольку программные модели, предполагающие, что указатели — это «фиксированные маркеры» в памяти, не сочетаются с аксессорами SYCL, существует расширение USM, делающее аксессоры ненужными. USM — это тяжкая ноша для программиста с точки зрения движения и синхронизации данных, но оно помогает сохранять совместимость с легаси кодом, использующим указатели.
Наша первая SYCL программа
Приведем только что описанную программу SYCL SAXPY целиком:
Полный исходный код примера имеется на GitHub.
Unified Shared Memory (USM)
Хотя приведенная выше программа полностью функциональна и может быть использована на множестве платформ, некоторые пользователи сочтут ее довольно громоздкой. Кроме того, она несовместима с библиотеками и фреймворками, управляющими памятью с помощью указателей. Чтобы решить эту проблему в SYCL 1.2.1, Intel разработал расширение в DPC++ под названием Unified Shared Memory (USM), поддерживающее управление памятью с помощью указателей.
USM поддерживает две важных модели применения. Первая — автоматическое перемещение данных между хостом и устройством в разделяемой памяти. Вторая — для принудительного перемещения данных в отдельную память устройства и из нее.
Более подробно это приводится в предварительной спецификации SYCL 2020, но все что нужно знать для начала, вы найдете тут. Аргумент
q — это очередь, ассоциированная с устройством, где размещенные данные будут находиться (временно или постоянно).
Если мы используем device allocation, данные должны быть перемещены явно, то есть с использованием SYCL метода
memcpy, который работает так же, как и std::memcpy (то есть, цель находится слева):
Если мы используем USM, то аксессоры не требуются, это означает, что мы можем упростить код ядра до:
Вы найдете полностью рабочие примеры обеих версий USM в репозитории под названиями saxpy-usm.cc и saxpy-usm2.cc соответственно.
Сжатый синтаксис SYCL 2020
И последнее — на случай, если вам до сих пор интересно, зачем в каждой из программ был нужен opaque handler
h, вроде как, совсем не нужный. Ниже приведена эквивалентная реализация, добавленная в предварительную спецификацию SYCL 2020. Более того, мы можем также воспользоваться лямбда-именами, они включены как опция в предварительную спецификацию SYCL 2020. Вместе эти два небольших изменения делают код ядра SYCL такого же размера, как и исходный цикл С++, показанный в самом начале:
Мы начали с трех строк кода, который работал последовательно на CPU и закончили тремя строками кода, который работает параллельно на CPU, GPU, FPGA и прочих устройствах. Очевидно, что не все будет таким простым, как SAXPY, но, по крайней мере, вы теперь знаете, что SYCL не усложняет простые вещи, и что он построен на основе современных фишек С++ и универсальных концепций типа «parallel for», а не на чем-то принципиально новом, требующем дополнительного изучения.
