SDAccel это система программирования на OpenCL для ПЛИС фирмы Xilinx. В настоящее время всё более обостряется проблема разработки проектов для ПЛИС на традиционных языках описания аппаратуры, таких как VHDL/Verilog. Одним из методов решения проблемы является применение языка C++. OpenCL это один из вариантов применения языка С++ для разработки прошивок ПЛИС.

Небольшое вступление про фазовые переходы

Мне пришлось заняться программированием ПЛИС в далёком 2000 году. В то время компьютеры были не очень мощными, а ПЛИС — совсем маленькими. Я работал c ПЛИС серии MAX7000 фирмы Altera. Для разработки прошивок использовалась замечательная система MaxPlus II. Основным средством там был графический редактор. Поддержка VHDL и Verilog уже была, но была очень слабой. Поддерживались только синтезируемое подмножество VHDL, Verilog. Зато можно было получить VHDL модель готовой ПЛИС с sdf файлом временных задержек. И соотношение мощности компьютера и объёма ПЛИС позволяло провести моделирование всего проекта ПЛИС с временными задержками. Сейчас об этом можно только мечтать. Примерно в это время начался фазовый переход в разработке проектов ПЛИС. Это был переход от схемного ввода к использованию VHDL/Verilog для моделирования отдельных узлов и всего проекта. В нашей компании он совпал с переходом от Altera и MaxPlus II к Xilinx и ISE. У нас этот переход завершился в 2004 году.

В данный момент идёт второй фазовый переход. Он связан с переходом разработки проектов ПЛИС от VHDL/Verilog к языку С++. Дело в том, что при современном соотношении мощности компьютера и объёма ПЛИС провести сеанс моделирования проекта ПЛИС на VHDL/Verilog практически невозможно. Сеанс моделирования может длиться от нескольких часов до нескольких дней. Такое время можно позволить для окончательной верификации проекта, но не для разработки.

Что такое OpenCL ?

Система OpenCL была предложена в 2008 году компанией Apple. В дальнейшем была организована ассоциация «Khronos Group» в которую вошли ведущие компании, такие как INTEL, NVIDIA, AMD, ARM, GOOGLE, SONY, SAMSUNG и много других. Кроме OpenCL там развиваются и другие системы, например OpenXR — система виртуальной реальности.

OpenCL — это система проектирования на основе С++ для гетерогенных систем таких как:

• обычные процессоры
• многопроцессорные кластеры
• графические процессоры
• ПЛИС

OpenCL определяет модель системы, расширения языка C++, библиотеку функций для HOST компьютера.

Большое время моделирования связано с моделированием на уровне тактовой частоты. Применение языка Си удаляет сигнал тактовой частоты из описания проекта. В проекте остаются только операции с данными. Это позволяет на несколько порядков увеличить скорость моделирования и разработки.

Одной из первых заметных систем программирования на Си является система Catapult компании Mentor Graphics. Эта система появилась в 2004 году и успешно используется в том числе компанией Microsoft для реализации своего поискового сервера Bing с использованием ПЛИС Altera.

Фирма Xilinx примерно в 2013 году выпустила Vivado HLS, которая позволяет разрабатывать отдельные компоненты на С++ и впоследствии включать их в основной проект. На основе Vivado HLS созданы ещё несколько продуктов:

• SDSoc — ускорение отдельных функций. Система предназначена только для Zynq (это микросхема в которой в одном корпусе есть ПЛИС и процессор АРМ). Система уже доступна.
• SDAccel — система программирования на OpenCL. Система доступна, но не всем.
• SDNet — система проектирования сетевых приложений. Пока не доступна и говорить о ней ещё рано.

SDSoc и SDAccel характерны тем, что проект ПЛИС уже отходит на второй план. На первом плане — алгоритм. Обе системы позволяют позволяют провести моделирование на уровне исходного алгоритма написанного на С/C++ и далее перевести его на ПЛИС. Это позволяет резко увеличить сложность алгоритма. И не случайно, что сейчас обе эти системы внедряются в обработку изображений.

Если сравнить программирование для ПЛИС на VHDL/Verilog и на С/С++, то напрашивается аналогия между программированием для обычных процессоров на С/C++ и на ассемблере. Да, на ассемблере можно сделать более компактный и быстрый код. Но на С/С++ можно написать более сложную программу.

Модель вычислителя

Система состоит из HOST компьютера и вычислителя, которые связаны между собой по шине. В большинстве случаев это шина PCI Express. Однако Altera уже предлагает решения для своих ПЛИС со встроенным процессором АРМ. В этом случае используется шина AXI. По некоторым слухам компания Intel (которая купила Altera) разрабатывает процессор Xeon со встроенной ПЛИС. Основной системой проектирования там будет OpenCL, а для взаимодействия между процессором и ПЛИС будет использоваться QPI.

Внутри вычислителя расположен один или несколько блоков «Compute Unit», каждый из которых состоит из одного или нескольких «Processing Element». На этом уровне есть принципиальная разница между графическими процессорами и ПЛИС. Если в графическом процессоре количество «Processing Element» определено (хотя оно разное в разных моделях), то в ПЛИС это может меняться в зависимости от задачи.

Стандарт определяет несколько классов памяти:

• HOST Memory — память доступная приложению на HOST компьютере. Обычно это оперативная память компьютера.
• Global Memory — память доступная для HOST и для вычислителя. Обычно это динамическая память подключённая к ПЛИС или к графическому процессору.
• Global Constant Memory — память доступная по чтению и записи для HOST и только для чтения на вычислителе.
• Local Memory — память доступная только в пределах одного «Compute Unit»
• Private Memory — память доступная только в пределах одного «Processing element»

Дополнительно Xilinx вводит «Global OnChip Memory» — память доступная всем «Compute Unit».
Упрощённый алгоритм работы:

1. HOST проводит инициализацию устройства
2. HOST загружает программу в вычислитель
3. HOST подготавливает данные в HOST Memory
4. HOST запускает DMA канал для передачи данных из HOST Memory в Global Memory и ожидает завершение DMA
5. HOST запускает на выполнение вычислитель и ожидает завершение вычисления.
6. HOST запускает DMA канал для передачи результата из Global Memory в HOST Memory и ожидает завершение DMA.
7. HOST использует результаты вычисления.

Важно отметить следующее — всё общение между HOST и вычислителем идёт через Global Memory. В более сложных алгоритмах можно параллельно с вычислениями передавать данные для следующего цикла.

Что такое kernel ?

Kernel — это базовое понятие OpenCL. Собственно говоря, это функция которая выполняется на одном «Processing Element». Несколько kernel могут выполняться в рамках одного «Compute Unit». Это основной способ обеспечения параллельности операций для графических процессоров.
Пример определения функции:

__kernel void krnl_vadd(
     __global int* a,
     __global int* b,
     __global int* c,
     const int length);

В отличие от обычного описания здесь появляются новые ключевые слова, они как раз определены в стандарте OpenCL.

• __kernel — определяет функцию которая будет выполняться на вычислителе.
• __global — определяет что данные расположены в глобальной памяти.

SDAccel предлагает три способа реализации kernel:

• стандарт OpenCL
• C++ — при этом будут использованы все возможности Vivado HLS
• VHDL/Verilog — при этом будут использованы все возможности ПЛИС

Главное различие в реализации для GPU и ПЛИС

На примере простой функции сложения двух векторов очень удобно проследить главное различие в эффективной реализации кода для графических процессоров и для ПЛИС.

Функция сложения для GPU будет выглядеть так:

__kernel void 
krnl_vadd(
     __global int* a,
     __global int* b,
     __global int* c,
     const int length) 
{

  int idx = get_global_id(0);	

  c[idx] = a[idx] + b[idx];

  return;
}

А для ПЛИС вот так:

__kernel void __attribute__ ((reqd_work_group_size(1, 1, 1)))
krnl_vadd(
     __global int* a,
     __global int* b,
     __global int* c,
     const int length) {


  for(int i = 0; i < length; i++){
    c[i] = a[i] + b[i];
  }

  return;
}

Обратите внимание — в версии для GPU не используется параметр length. Предполагается, что для каждого элемента вектора будет запущен свой экземпляр kernel. Каждый экземпляр получит свой номер idx и выполнит сложение. Количество одновременно запущенных экземпляров будет определяться возможностями данного GPU. Если вектор будет слишком большим, то будет несколько запусков. Для ПЛИС так тоже можно сделать, но это не очень эффективно. Лучшие результаты даёт вариант в котором используется только один «Compute Unit» и один «Processing Elenet». Обратите внимание — перед объявлением функции добавился атрибут reqd_work_group_size( 1, 1, 1), а внутри самой функции есть цикл. Значение атрибута 1,1,1 означает что будет использоваться только один kernel. И это знание будет использоваться для оптимизации вычислительной структуры. Сам цикл с помощью дополнительных атрибутов может быть развёрнут в параллельную вычислительную структуру. Наилучший результат будет достигнут если length будет константой.

SDAccel

Начиная с версии 2016.3 SDAccel и SDSoc объединены в один пакет под названием SDx. SDSoc работает в Windows и Linux. SDAccel работает только под некоторыми версиями Linux, в частности — CentOs 6.8; Разумных объяснений таким ограничениям нет, надеюсь в будущем SDAccel будет работать и под Windows. Пакет SDx сделан на основе Eclipse. В нём добавляется тип проекта «Xilinx SDx». При создании проекта требуется выбрать платформу. Пока выбор небольшой. На рисунке представлен вид окна выбора платформы:



Платформа определят модуль и базовую прошивку ПЛИС. В SDAccel используется технология частичной перезагрузки (Partial Reconfiguration). Требуется соответствие между базовой прошивкой, которая загружена в ПЛИС и той на основе которой формируется проект SDAccel. Это соответствие поддерживается названием и версией платформы. Обратите внимание, верхняя строчка — это модуль FMC126P. Я пытаюсь создать для него платформу, пока неудачно.

Ещё один важный скришот — свойства проекта:



Обратите внимание на поле «HW Functions».

• binary_container_1 — это то, что будет загружено в ПЛИС.
• kernel-vadd — это имя функции
• Колонка «Comput Units» — по сути, это количество параллельных реализаций функции
• Колонка «Max Memory Ports» — разрешение дополнительной оптимизации при обращении к глобальной памяти

Очень важным является правое верхнее поле: «Active build configuration». Собственно говоря здесь заключена вся суть этой системы. Возможно три варианта:

• Emulation-CPU — реализация OpenCL на процессоре
• Emulation-HW — реализация OpenCL на симуляторе Vivado
• System — реализация OpenCL на выбранной аппаратной платформе

Результатом компиляции будет выполняемый файл, он кстати имеет расширение .exe, и файл с расширением .xclbin; Это binary_container с реализацией функций kernel.

Для трёх вариантов выполнения формируется разная среда выполнения OpenCL. Вариант Emulation-CPU самый быстрый для запуска. Компиляция и запуск производятся очень быстро. Именно в этом режиме надо проверять алгоритм.

Вариант Emulation-HW более долгий для компиляции и выполнения. В этом режиме вызывается Vivado HLS, производится синтез кода для VHDL/Verilog/SystemC и запускается симулятор Vivado для выполнения kernel. По резудьтатам компиляции можно определить занимаемые ресурсы и оценить задержки на выполнение. Моделирование может быть долгим, поскольку здесь уже есть тактовая частота и мы получаем все связанные с этим проблемы. Хотя наверняка для PCI Express и SODIMM используются упрощённые модели, что увеличивает скорость моделирования.

Вариант System является рабочим. Компиляция включает в себя трассировку ПЛИС, что является достаточно долгим процессом. Небольшой проект для ADM-PCIE-KU3 разводится около часа. Для запуска требуется установить драйвер устройства, который поставляется вместе с платформой. При запуске производится загрузка binary_container в ПЛИС с использованием технологии Partial Reconfiguration. Сама загрузка тоже не быстрая, около минуты. С чем это связано я объяснить не могу.

Программа для HOST

Стандарт OpenCL определяет API. На сайте Khronos Group все функции хорошо описаны. Но выглядит всё это достаточно мрачно. Однако Xilinx здесь тоже упростил нам жизнь. В состав примера vector_addition входят файлы xcl.h и xcl.cpp, в которых описаны самые необходимые функции для работы с одним устройством. Вот они:

• xcl_world_single(), xcl_world_release() — инициализация и завершение работы с устройством
• xcl_malloc() — выделение буфера в глобальной памяти на устройстве
• xcl_import_binary() — загрузка binary_container
• xcl_set_kernel_arg() — установка аргументов для функции kernel
• xcl_memcpy_to_device() — передача данных на устройство
• xcl_memcpy_from_devce() — передача данных из устройства
• xcl_run_kernel3d() — запуск функции на выполнение

Конечно программа для HOST может быть не одна. Вполне возможно сделать отдельный проект и подключить какую-либо систему Unit тестирования, например Google Test, для проверки реализации функций на ПЛИС.

А что внутри ПЛИС ?

В каталоге компонентов есть такой симпатичный элемент «SDAccel OpenCL Programmable Region»



Вот именно в него и будет загружен binary_container. Видно, что элемент имеет крайне малое количество связей. Есть шина S_AXI для управления, шина M_AXI для доступа к глобальной памяти, ну и сигналы тактовой частоты и сброса. Предполагается, что в ПЛИС есть узел DMA, контроллер динамической памяти и центральный узел axi_interconnect.

Блок SDAccel можно раскрыть, внутри он будет выглядеть так:



Не очень хорошо, но видно что есть два блока axi_interconnect, а между ними четыре блока kernel. Из такой структуры следует рекомендация не использовать большое количество kernel, поскольку на каждый блок потребуется своя шина AXI. Не рекомендуется использовать более 16 шин.

Потенциальные преимущества и реальные недостатки

Главным преимуществом системы является возможность реализации сложных алгоритмов для работы с большими массивами данных. Конечно понятия «сложный алгоритм» и «большой массив» являются условными. По моему субъективному мнению, применение системы будет эффективным для тех алгоритмов, для проверки которых требуется более 1 Мбайта тестовых данных. В первую очередь это конечно алгоритмы обработки изображений.

Другим потенциальным преимуществом является возможность перехода на другую аппаратуру. Например с ПЛИС Xilinx на ПЛИС Altera.

Основные недостатки:

• Это новая система, наверняка ещё есть необъяснимые баги
• Работа только под ограниченное количество вариантов Linux. Под Windows — не работает.
• Под вопросом эффективность трансляции с языка С++ на VHDL/Verilog. Хотя существует возможность реализации kernel на VHDL/Verilog.

Первое знакомство состоялось, что дальше

При дальнейшем изучении SDAccel я планирую следующее:

1. Изучение эффективных методов работы с памятью, измерение скорости работы
2. Разработка платформы для модуля FMC126P
3. Реализация узла свёртки на основе библиотеки FPFFTK от Александра Капитанова ( capitanov )

P.S. Кстати, OpenCL не поддерживает <stdio.h>, однако printf там есть. В том числе printf работает и при реализации на ПЛИС.