Всем привет!
В прошлой статье я запустил простой OpenCL пример на FPGA фирмы Altera:
// ACL kernel for adding two input vectors
__kernel void vector_add( __global const uint *restrict x,
__global const uint *restrict y,
__global uint *restrict z )
{
// get index of the work item
int index = get_global_id(0);
// add the vector elements
z[index] = x[index] + y[index];
}
Я намеренно не углублялся в детали и показал верхушку айсберга: процесс разработки, сборку проекта, запуск на системе.
При подготовке первой статьи мне стало дико интересно, во что превращаются (со стороны FPGA) эти строчки. Понимание архитектуры даст возможность что-то соптимизировать и понять на что уходят ресурсы, а так же что хорошо и плохо для этой системы.
В этой статье мы попробуем вскрыть ядро и найти ответы на следующие вопросы:
- Какая у него архитектура?
- Как происходит его настройка? Как попадают данные на обработку?
- На какой частоте он работает? Чем это определяется?
- Можно ли просимулировать только ядро в RTL-симуляторах?
- Какие блоки занимают больше всего ресурсов? Можно ли как-то это соптимизировать?
Давайте взглянём на его внутренности! Добро пожаловать под кат!
Как это видит Альтера
Перед тем, как пристально изучать проект для FPGA, обратимся к различным презентациям от вендора: что они рассказывают про реализацию на высокоуровневом (маркетинговом) языке.
Рекомендую глянуть большую презентацию-введение в OpenCL от Альтеры
Harnessing the Power of FPGAs using Altera’s OpenCL Compiler (осторожно, больше ста слайдов, ~16 МБ).
Прошивка состоит из:
- IP-ядер, которые обеспечивают доступ до периферии (PCIe, внешняя память (DDR, QDR)).
- Ядер, реализованые по принципу конвейера. В них происходят вычисления, описанные в OpenCL ядрах.
- Инфраструктура: Global и Local Memory Interconnect.
Интерконнект — это деление общей шины между модулями, которые являются мастерами и слейвами (ведущими и ведомыми).
В нашем случае мастера — это ядра, которые читают/пишут данные как в глобальную память (это может быть как память хоста, так и внешняя память), так и в локальную (внутреннюю), которую можно назвать кэшом. В результате процесса арбитража и мультиплексирования данных появляются модули, которые, как мы увидим ниже, могут отъедать значительное количество ресурсов.
Для удобства протокол общения между модулями стандартизируют. Altera в своих проектах использует интерфейсы типа Avalon: Avalon-MM (Memory Mapped) и Avalon-ST (Streaming). Я на этом подробно останавливаться не буду: читатель может самостоятельно про это почитать тут. В этой статье большинство интерконнекта будет именно интерфейса Avalon-MM.
Еще раз сделаю акцент на том, что всё это получается автоматически из описания ядра на OpenCL.
Результаты после обновления
В прошлой статье я описывал результаты сборки исходя из работы на версии Quartus 14.1.
Не так давно вышла версия 15.1, и я решил посмотреть, есть ли там большие различия. Для этого я перегенерировал исходники и пересобрал их новой версией.
Увы, в визуализаторе и профилировщике OpenCL никаких изменений не произошло (с виду): их вид всё еще оставляет желать лучшего.
Отчет о сборке с --profile (с профилирующими счетчиками):
+-----------------------------------------------------------------------------------+
; Fitter Summary ;
+---------------------------------+-------------------------------------------------+
; Fitter Status ; Successful - Sun Nov 22 13:18:14 2015 ;
; Quartus Prime Version ; 15.1.0 Build 185 10/21/2015 SJ Standard Edition ;
; Family ; Cyclone V ;
; Device ; 5CSEMA5F31C6 ;
; Timing Models ; Final ;
; Logic utilization (in ALMs) ; 5,472 / 32,070 ( 17 % ) ;
; Total registers ; 10409 ;
; Total pins ; 103 / 457 ( 23 % ) ;
; Total block memory bits ; 127,344 / 4,065,280 ( 3 % ) ;
; Total RAM Blocks ; 44 / 397 ( 11 % ) ;
; Total PLLs ; 2 / 6 ( 33 % ) ;
; Total DLLs ; 1 / 4 ( 25 % ) ;
+---------------------------------+-------------------------------------------------+
По сравнению с предыдущей версией компилятора проект похудел примерно на 100 ALM.
А вот отчет сборки без профилирующих счетчиков:
+-----------------------------------------------------------------------------------+
; Fitter Summary ;
+---------------------------------+-------------------------------------------------+
; Fitter Status ; Successful - Sun Nov 22 13:51:21 2015 ;
; Quartus Prime Version ; 15.1.0 Build 185 10/21/2015 SJ Standard Edition ;
; Family ; Cyclone V ;
; Device ; 5CSEMA5F31C6 ;
; Timing Models ; Final ;
; Logic utilization (in ALMs) ; 4,552 / 32,070 ( 14 % ) ;
; Total registers ; 7991 ;
; Total pins ; 103 / 457 ( 23 % ) ;
; Total block memory bits ; 127,344 / 4,065,280 ( 3 % ) ;
; Total RAM Blocks ; 44 / 397 ( 11 % ) ;
; Total PLLs ; 2 / 6 ( 33 % ) ;
; Total DLLs ; 1 / 4 ( 25 % ) ;
+---------------------------------+-------------------------------------------------+
Как видим, около 1000 ALM занимают профилирующие счетчики и логика, которые их «вычитывает».
В дальнейшем именно этот отчет мы будем использовать для анализа, что сколько занимает.
Первый взгляд на проект
Напомню, что проект выложен на на гитхабе.
Файл проекта называется незамысловато: top.qpf (QPF — Quartus Project File), самый главный модуль top.v, который по факту содержит экземпляр модуля system и простой счетчик, который отображается на светодиодах.
system (4535 ALM)
system — это автосгенеренный модуль с помощью Qsys. Qsys — это GUI-ишная тулза, которая позволяет соединять различные IP-блоки, автоматически генерируя код модулей, которые необходимы для интерконнекта, перехода с одной частоту на другую и пр.
Модули:
- vector_add_system (2141 ALM) — это модуль, который реализует то, что мы написали в ядре vector_add.
- acl_iface (2343 ALM) — инфраструктура, которая обеспечивает более удобный доступ и взаимодейстие с ядром.
Интерфейсы:
- avs_vector_add_cra — Avalon-MM для управления ядром.
- avm_memgmem0_port_0_0_rw — Avalon-MM для доступа к DDR памяти. Ширина данных — 256 бит.
acl_iface (2343 ALM)
Модули:
- pll (0 ALM): PLL, которая получает клок pll_outclk0 (100 МГц) из config_clk (50 МГц, поступает с внешнего генератора).
- acl_kernel_clk (1057 ALM): еще одна PLL: она генерирует клок, который подается на ядро. Она имеет интересный нюанс: о ней более подробно поговорим позже.
- acl_kernel_interface (439 ALM) : обеспечивает «взаимодействие» ядра и процессора (через интерфейс управления и прерывание).
- clock_cross_kernel_mem1 (82 ALM): он занимается «согласованием» интерфейсов, которые работают на разных частотах (CDC ).
- hps (0 ALM): это инстанс HPS (Hard Processor System). Никакой логики в FPGA он не занимает, т.к. это аппаратное ядро.
Интерфейсы:
- f2h_sdram0 — Avalon-MM интерфейс для доступа к DDR памяти. Ширина данных — 256 бит, а частота работы — pll_outclk0 (100 МГц).
- h2f_lw — AXI интерфейс. Через него CPU (ARM) имеет возможность управлять и настраивать систему используя контрольные/статусные регистры кернела и т.д.
Если сложим суммарную ёмкость этих модулей, то сумма не сойдется. Дело в том, что Qsys по умолчанию не показывает модули типа interconnect. Для отображения их необходимо нажать Show System With Qsys Interconnect в меню System. После этого можно увидеть, что есть модули вида mm_interconnect_*, которые занимают 568 и 195 ALM.
vector_add_system (2141 ALM)
Архитектуру этого модуля нельзя посмотреть в GUI: для понимания как он работает погружаемся в Verilog.
Примерная схема выглядит так:
- vector_add_system_interconnect_* (443 ALM) — модули интерконнекта, которые проводят арбитраж и мультиплексирование интерфейса avm_memgmem0_port_0_0_rw
- LSU_X (235), LSU_Y (239) — вычитывают данные из глобальной памяти для векторов (аргументы ядра x и y соответственно).
- LSU_Z (424 ALM) — записывает результат вычислений в глобальную память (аргумент z).
- acl_id_iterator (228 ALM), acl_work_group_dispatcher (149 ALM) — они выдают задание для выполнения ядру (показывают, какой элемент надо обработать).
- acl_kernel_finish_detector (144 ALM) — определяет, когда ядро закончило работу.
Примечание:
LSU-модули являются инстансами одного модуля (lsu_top) и имеют названия lsu_local_bb0_ld_, lsu_local_bb0_ld__u0 и lsu_local_bb0_st_add. Для удобства я придал им более «человечные» названия. Более подробно про LSU мы поговорим ниже.
Как работает ядро:
- Происходит настройка через CRA, запускается обработка.
- LSU_X и LSU_Y получают «команды» на чтение данных и делают запросы к глобальной памяти.
- Прочитанные данные буферизируются в памяти (FIFO), до тех пор, пока не будут готовы данные с обоих LSU.
- Как только данные есть в обоих FIFO, они отправляются на конвеер, который производит сложение.
- Результат попадает в LSU_Z, где они дожидаются возможности быть записанными в глобальную память.
- Как только обработано желаемое количество элементов, и все результаты записаны в память (нет отложенных записей) срабатывает kernel_finish_detector — выставляется прерывание kernel_irq.
Важно отметить, что три LSU будут между собой бороться за один интерфейс доступа к глобальной памяти — они являются мастерами интерфейса Avalon-MM.
Конвеер, который я обозначил в схеме как add_pipeline на самом деле в отдельный модуль не помещается: он просто расположен в файле vector_add.v в модуле vector_add_basic_block_0.
Сама строчка, которая осуществляет сложение двух 32-битных чисел, выглядит вот так:
assign local_bb0_add = (rstag_3to3_bb0_ld__u0 + rstag_3to3_bb0_ld_);
Логические элементы, которые будут созданы из этой строчки и делают всю полезную работу.
Всё остальное — инфраструктура, которая подгоняет данные к этой логике.
LSU (Load Store Unit)
Самым интересным модулем этого ядра является LSU. Давайте посмотрим, как он работает.
lsu_top — это по факту обертка над другими lsu_*-модулями, которые выбираются в зависимости от параметров READ и STYLE.
Из всех разновидностей у нас будет только две:
- LSU_READ_STREAMING — LSU_X, LSU_Y (READ = 1, STYLE = «STREAMING»)
- LSU_WRITE_STREAMING — LSU_Z (READ = 0, STYLE = «STREAMING»)
LSU_READ_STREAMING
Обратим внимание на параметры модуля:
BURSTCOUNT_WIDTH = 5;
MEMORY_SIDE_MEM_LATENCY = 89;
BURSTCOUNT_WIDTH показывает ширину сигнала avm_burstcount — при запросе по интерфейсу Avalon-MM там располагается количество слов, которое необходимо прочитать при бёрстной транзакции.
Если ширина сигнала равна пяти, то максимальное значение бёрста равно 16. Это явно следует из спецификации:
The value of the maximum burstcount parameter must be a power of 2.
A burstcount interface of width n can encode a max burst of size 2^(n-1).
For example, a 4-bit burstcount signal can support a maximum burst count of 8.
The minimum burstcount is 1.
Это значит, что максимум за один запрос будет прочитано 16 256-битных слов — это 4096 Кбит или 128 32-битных чисел (мы складываем именно 32-битные целые числа).
MEMORY_SIDE_MEM_LATENCY влияет на количество слов FIFO в lsu_burst_read_master. Эта FIFO служит для буферизации читаемых данных с глобальной памяти.
Как определяется количество слов для неё:
localparam MAXBURSTCOUNT=2**(BURSTCOUNT_WIDTH-1);
// Parameterize the FIFO depth based on the "drain" rate of the return FIFO
// In the worst case you need memory latency + burstcount, but if the kernel
// is slow to pull data out we can overlap the next burst with that. Also
// since you can't backpressure responses, you need at least a full burst
// of space.
// Note the burst_read_master requires a fifo depth >= MAXBURSTCOUNT + 5. This
// hardcoded 5 latency could result in half the bandwidth when burst and
// latency is small, hence double it so we can double buffer.
localparam _FIFO_DEPTH = MAXBURSTCOUNT + 10 + ((MEMORY_SIDE_MEM_LATENCY * WIDTH_BYTES + MWIDTH_BYTES - 1) / MWIDTH_BYTES);
// This fifo doesn't affect the pipeline, round to power of 2
localparam FIFO_DEPTH = 2**$clog2(_FIFO_DEPTH);
_FIFO_DEPTH = 16 + 10 + ((89 * 4 + 32 - 1)/32) = 39
Округляем вверх до числа, кратного степени двойки:
FIFO_DEPTH = 64
Вывод:
будет выделен буфер (кэш) на 64 слова по 256 бит.
На самом для фиксирования этого факта не обязательно было ковырять исходники: достаточно глянуть секцию RAM Summary отчета от сборке. Наш расчет оказался верным, причем в отчете видно, что будет использовано 7 блоков M10K. Семь блоков это 10240 бит * 7 = 70 Кбит, вместо ожидаемых 256 бит * 64 = 16 Кбит.
Почему так произошло?
В FPGA внутренняя память — это много маленьких блоков, которые могут быть по разному настроены.
Посмотреть как можно сконфигурировать блок M10K (а именно он составляет основу в чипах семейства Cyclone V) можно тут.
Максимальная длина слова в блоке памяти — 40 бит, если необходимо создать слово в 256 бит, то надо 256/40 = 6.4 -> 7 блоков, которые и получились. Из-за того, что количество слов в памяти выбрано 64, то каждый блок будет сконфигурирован как 64x40, и оставшиеся 75% памяти будут просто не использоваться.
На что влияет размер бёрста и размер кэша?
- Чем больше бёрст, тем больше мы можем прочитать за один запрос, но при этом будут блокироваться остальные запросы к памяти (у нас три мастера, которые хотят общаться с внешней памятью).
- Чем больше кэш, тем больше данных есть в «запасе» для обработки, пока читается новая порция данных. Минусов большего кэша я не знаю, кроме расхода ресурсов. В данном случае можно было сделать кэш с количеством слов равное 256 и было бы потрачено такое же количество блоков M10K.
LSU_WRITE_STREAMING
Поступающие 32 битные данные (результат сложения) кладутся по очереди в свои FIFO. Как только набирается в каждой из них по MAXBURSTCOUNT (для этого модуля этот параметр тоже равен 16), то происходит транзакция записи в память. Каждая из таких FIFO имеет ширину данных размером 32. Таких фифошек FIFO восемь штук (256/32).
На какое количество данных рассчитаны эти фифошки?
Расчет возьмем из кода lsu_streaming_write. Для этого модуля параметр. MEMORY_SIDE_MEM_LATENCY равен 32.
localparam MAXBURSTCOUNT=2**(BURSTCOUNT_WIDTH-1);
localparam __FIFO_DEPTH=2*MAXBURSTCOUNT + (MEMORY_SIDE_MEM_LATENCY * WIDTH + MWIDTH - 1) / MWIDTH;
localparam _FIFO_DEPTH= ( __FIFO_DEPTH > MAXBURSTCOUNT+4 ) ? __FIFO_DEPTH : MAXBURSTCOUNT+5;
// This fifo doesn't affect the pipeline, round to power of 2
localparam FIFO_DEPTH= 2**($clog2(_FIFO_DEPTH));
MAXBURSTCOUNT = 2^4 = 16
__FIFO_DEPTH = 2 * 16 + ( 32 * 32 + 256 - 1)/256 = 36 + 5 = 41
_FIFO_DEPTH = 41
Округляем вверх до кратного степени двойки:
FIFO_DEPTH = 64
Подтвержаем отчетом: 64 * 32 = 2048 бит (1 M10K).
Так как FIFO полностью отдельные, то на каждую FIFO выделятся по одному блоку M10K, что приводит к 8 блокам M10K, против 7 блоков M10K в lsu_read_streaming.
Почему сделали 8 FIFO, хотя можно было сделать одну, но широкую? Скорее всего так проще сделать (не надо отдельно хранить количество валидных слов).
Как вычисляются параметры LSU?
Попробуем разобраться, откуда такие числа возникли:
Есть подозрение, что эти настройки беруться из файла, который описывает плату (altera/15.1/hld/board/de1soc/de1soc_sharedonly/board_spec.xml).
Находим строчку, которая связана с глобальной памятью:
<!-- One DDR3-800 DIMM, 256-bit data -->
<global_mem max_bandwidth="6400">
<interface name="acl_iface" port="kernel_mem0" type="slave" width="256" maxburst="16" latency="240" address="0x00000000" size="0x40000000"/>
</global_mem>
Для разъяснений этих параметров обратимся к Altera SDK for OpenCL: Custom Platform Toolkit User Guide глава XML Elements, Attributes, and Parameters in the board_spec.xml File.
max_bandwidth — The maximum bandwidth of all global memory interfaces combined in their current configuration. The Altera Offline Compiler uses max_bandwidth to choose an architecture suitable for the application and the board. Compute this bandwidth value from datasheets of your memories.
Перевод
max_bandwidth — Максимальная пропускная способность всех интерфейсов для глобальной памяти. Altera Offline Compiler использует max_bandwidth для выбора архитектуры, которая лучше всего подходит для конкретной платы и приложения. Вычислите эти значения исходя из параметров используемой памяти.
К сожалению нет пояснений в каких единицах и как это считается: с одной стороны в профилировщике писалось 6400 MB/s, с другой стороны по расчетам 6400 MB/s никак не получаются: 400 (МГц, тактовая частота DDR) * 32 (бит, ширина сигнала данных на DDR-интерфейсе) * 2 (работа по двум фронтам) = 25600 Mb/s = 3200 MB/s. Либо надо считать в обе стороны?
max_burst — Maximum burst size for the slave interface.
Перевод
max_burst — Максимальный размер берста для ведомого (слейв) интерфейса.
В нашем случае — 16, что дает BURSTCOUNT_WIDTH = 5. Но почему именно 16? Интерфейс fpga2hps_sdram поддерживает max_burstcount = 128. 16 — это какое-то магическое число, подходит всем? :)
latency — An integer specifying the time in nanoseconds (ns) for the memory interface to respond to a request. The latency is the round-trip time from the kernel issuing the board system a memory read request to the memory data returning to the kernel. For example, the Altera DDR3 memory controller running at 200 MHz with clock-crossing bridges has a latency of approximately 240 ns.
Перевод
latency — Целое число, которые показывает время в наносекундах, необходимое
интерфейсу памяти для ответа. Задержка — это время от запроса на чтение, до получения данных в ядре. Например, Altera DDR3 контроллер, работающий на частоте 200 МГц в связке с модулем для перехода на другую частоту, имеет задержку около 240 нс.
интерфейсу памяти для ответа. Задержка — это время от запроса на чтение, до получения данных в ядре. Например, Altera DDR3 контроллер, работающий на частоте 200 МГц в связке с модулем для перехода на другую частоту, имеет задержку около 240 нс.
Допустим, что в нашем случае тоже задержка 240 ns. Очевидно, что размерность MEMORY_SIDE_MEM_LATENCY это количество тактов (да и комментарий это подсказывает: Latency in cycles between LSU and memory).
Проведём несколько экспериментов, изменяя значения в board_spec.xml (maxburst, latency) и структуру ядра (количество аргументов, которые складываются (readers)). Следим за значением параметра MEMORY_SIDE_MEM_LATENCY у обоих модулей (LSU_X (lsu_read_streaming) и LSU_Z (lsu_write_streaming) ).
|--------------------------------------------------------|
| maxburst | latency | readers | MEMORY_SIDE_MEM_LATENCY |
| | | |-------------------------|
| | | | LSU_X | LSU_Z |
|--------------------------------------------------------|
| 16 | 0 | 1 | 25 | 16 |
| 16 | 100 | 1 | 45 | 16 |
| 16 | 240 | 1 | 73 | 16 |
|--------------------------------------------------------|
| 16 | 0 | 2 | 41 | 32 |
| 16 | 100 | 2 | 61 | 32 |
| 16 | 240 | 2 | 89 | 32 |
|--------------------------------------------------------|
| 16 | 0 | 3 | 57 | 48 |
| 16 | 100 | 3 | 77 | 48 |
| 16 | 240 | 3 | 105 | 48 |
|--------------------------------------------------------|
| 32 | 0 | 1 | 41 | 32 |
| 32 | 100 | 1 | 61 | 32 |
| 32 | 240 | 1 | 89 | 32 |
|--------------------------------------------------------|
| 32 | 0 | 2 | 73 | 64 |
| 32 | 100 | 2 | 93 | 64 |
| 32 | 240 | 2 | 121 | 64 |
|--------------------------------------------------------|
| 32 | 0 | 3 | 105 | 96 |
| 32 | 100 | 3 | 125 | 96 |
| 32 | 240 | 3 | 153 | 96 |
|--------------------------------------------------------|
Какие прослеживаются зависимости:
- При увеличении latency и фиксировании maxburst и readers LSU_X_MEMORY_SIDE_MEM_LATENCY возрастает на значение равное latency/5. Скорее всего 5 — это 5 ns (нас отсылают к магической частоте 200 МГц?).
- При увеличении количества читателей LSU_X_MEMORY_SIDE_MEM_LATENCY возрастает на значение бёрста.
- LSU_Z_MEMORY_SIDE_MEM_LATENCY линейно зависит от количества элементов, которые хотят получать доступ к глобальной памяти и от значения максимально бёрста.
Виднеются формулы:
- LSU_X_MEMORY_SIDE_MEM_LATENCY = 9 + readers * maxburst + latency/5. (9 — это либо какое-то магическое число, либо еще одна характеристика ядра, до которой я не докопался. Возможно, это общая задержка ядра).
- LSU_Z_MEMORY_SIDE_MEM_LATENCY = maxburst * readers.
Примечание:
эти формулы только для конкретного ядра (реализации), для другого всё может быть иначе.
На какой частоте работает ядро
Генерацией тактового сигнала для ядра занимается модуль acl_kernel_clk.
В его основе лежит PLL, которая может динамически реконфигурироваться (менять выходную частоту).
Если откроем этот модуль в Qsys или system_acl_iface_acl_kernel_clk_kernel_pll.v, то увидим, что эта PLL генерирует два сигнала — 140 МГц (kernel_clk) и 280 МГц (kernel_clk2x). Сразу скажу, что kernel_clk2x нигде не используется.
Означает ли, что ядро всегда (и любое) будет работать только на частоте 140 МГц и его никак нельзя разогнать? Конечно же, нет.
140 МГц — это настройка для конкретно этой платы.
В завимости от того, какие логические элементы используются и как они соединены, то значение тактовой частоты, на которой схема будет гарантировано работать без сбоев может быть различным. Я затрагивал этот вопрос в статье про конвейеризацию.
Задача компилятора — расположить примитивы (логические элементы, блоки памяти и пр.) так, чтобы удовлетворить заданному требованию частоты. Это значит что:
- он не старается найти такое расположение, которое даст самую максимальную тактовую частоту.
- если в течении некоторого времени он, перебирая расположение элементов в чипе, понимает, что он развести не может, то он оставляет один из лучших вариантов (который был за время поиска).
Допустим вместо 140 МГц Quartus показывает максимальную тактовую частоту 135 МГц. Это значит, что:
- компилятор гарантирует, что если подать 135 МГц, то вычисления произойдут корректно, ничего не зависнет и пр. (если нет алгоритимических ошибок в самом коде, разумеется).
- если подать 140 МГц, то может так произойти, что всё будет хорошо. А может быть и нет. Это зависит от чипа — чипы с одной маркировкой могут немного отличаться, поэтому компилятор перестраховывается и расчитывает по худшему случаю.
Чаще всего после пересборки проекта FPGA разработчики смотрят отчет о сборке и интересуются: уложилось ли по частоте схема. Мы же в прошлой статье просто взяли бинарник и зашили его. Что будет, если компилятор не уложился в эти 140 МГц? Расчеты будут неверны?
Для того, чтобы скрыть от разработчиков эту проблему, Altera сделала очень интересную фишку (наверно, самую интересную из того, что я раскопал, когда игрался с Altera OpenCL SDK):
- После окончания сборки вызывается скрипт adjust_plls.tcl. Он получает максимальную допустимую частоту для ядра (Fmax), и генерирует файлы (pll_rom.mif и pll_rom.hex), которые используются для инициализации ROM в модуле pll_rom.
- Когда загружается FPGA, на логику подается заданная частота (140 МГц). Перед запуском ядра, вычитываются данные из ROM, и используя эти коэффициенты происходит перестройка PLL (через интерфейс реконфигурации). Как только реконфигурация закончилась на кернел уже подается нужная частота.
Скрытый текст
Судя по коду tcl-скрипта в ROM кладутся коэфициенты m, n, k, c0, c1, которые должны соответствовать тем, которые описаны в документации про перестройку Altera PLL.
Итого:
- На ядро будет подана та, частота которую можно подавать. Если логика оказалась слишком ёмкой, и не удалось уложиться в заданное число, то вычисления не сломаются — просто они будут медленее идти.
- Если же допустимая частота выше, то PLL будет настроено на это значение (вычисления ускорятся). Поиска расположения, которое даст максимальную частоту, не будет. Если есть ощущение, что еще можно «разогнать», то лучше вручную поднимать планку у частоты PLL.
Немного упрощаем сборку
Перед тем как мы продолжии познавать как устроено и настраивается ядро я сделаю небольшое отступление, которое может помочь вам, если захотите внести какие-то изменения (в ядро) или отлаживаться на железе.
Напомню процесс разработки: файл vector_add.aocx, который содержит прошивку FPGA получается из vector_add.cl.
Проблема заключается в том, что если вы внесли какие-то изменения в проект Квартуса, то они не попадут в *.aocx, т.к. при перезапуске утилиты aoc происходит копирование «дефолтного проекта» и перегенерация Verilog IP. Тем самым ваши изменения пропадут.
Утилита aoc является бинарником, но можно проследить, что при вызове:
$ aoc device/vector_add.cl -o bin/vector_add.aocx --board de1soc_sharedonly --profile -v
Происходит запуск скрипта на перле aoc.pl, который уже и делает всю полезную работу.
Можно напрямую вызывать этот скрипт, без использования утилиты aoc.
$ /home/ish/altera/15.1/quartus/linux64/perl/bin/perl /home/ish/altera/15.1/hld/share/lib/perl/acl/aoc.pl device/vector_add.cl --board de1soc_sharedonly --profile -v
Хорошо, что скрипт написан на интерпретируемом языке, а значит мы сможем разобраться, что он делает и внести свои изменения.
В самом начале скрипта описаны различные переменные, которые настраиваются через ключи (в том числе скрытые от пользователя в хелпе).
Так, там обнаруживается ключик --quartus, который вызывает только сборку квартуса и упаковку необходимых частей в *.aocx файл. Никакой перегенерации проекта (исходников) при этом не будет.
Так же для дополнительного удобства можно вывести лог сборки на консоль. Для этого надо в качестве stdout и stderr указать пустые строчки в вызове функции mysystem_full:
$return_status = mysystem_full(
{'time' => 1, 'time-label' => 'Quartus compilation', 'stdout' => '', 'stderr' => ''},
$synthesize_cmd);
Теперь мы можем легко вносить любые изменения в проект (играться с оптимизациями, добавлять SignalTap) и просто вызывать пересборку только проекта для FPGA, а не всего ядра с вызовом clang'a и перегенерацией кода.
Для проверки этого я добавил SignalTap на интерфейсы (а так же добавил 15 секундый sleep после загрузки ядра и стартом вычислений, чтобы я успел подключиться с помощью дебаггера).
Как управляется ядро
Интерфейс avs_vector_add_cra служит для настройки кернела: по адресам регистров записываются данные.
К сожалению, в открытом доступе я не нашел карты регистров и что как надо настраивать, поэтому придется сделаем небольшое исследование.
Все регистры описываются в vector_add.v и имеют адекватные названия.
Они являются 64-битными: [31:0] обозначают нижние 32 бита, а [63:32] — старшие.
0x0 - status
0x1 - 0x4 - profile
0x5 - [31:0] - work_dim
0x5 - [63:32] - workgroup_size
0x6 - [31:0] - global_size[0]
0x6 - [63:32] - global_size[1]
0x7 - [31:0] - global_size[2]
0x7 - [63:32] - num_groups[0]
0x8 - [31:0] - num_groups[1]
0x8 - [63:32] - num_groups[2]
0x9 - [31:0] - local_size[0]
0x9 - [63:32] - local_size[1]
0xA - [31:0] - local_size[2]
0xA - [63:32] - global_offset[0]
0xB - [31:0] - global_offset[1]
0xB - [63:32] - global_offset[2]
0xC - [31:0] - kernel_arguments[31:0] - input_x[31:0]
0xC - [63:32] - kernel_arguments[63:32] - input_x[63:32]
0xD - [31:0] - kernel_arguments[95:64] - input_y[31:0]
0xD - [63:32] - kernel_arguments[127:96] - input_y[63:32]
0xE - [31:0] - kernel_arguments[159:128] - input_z[31:0]
0xE - [63:32] - kernel_arguments[191:160] - input_z[63:32]
Исходя из названий, можно попытаться наугад что-то настроить и запустить, но давайте не рисковать, а просто узнаем что и в каком порядке туда пишется.
Запишем все транзакции по этому интерфейсу (с помощью SignalTap'a):
----------------------------------------------
| addr | write_data | byte_enable |
----------------------------------------------
| 0x5 | 0x00000000 0x00000001 | 0x0F |
| 0x5 | 0x000F4240 0x00000000 | 0xF0 |
----------------------------------------------
| 0x6 | 0x00000000 0x000F4240 | 0x0F |
| 0x6 | 0x00000001 0x00000000 | 0xF0 |
----------------------------------------------
| 0x7 | 0x00000000 0x00000001 | 0x0F |
| 0x7 | 0x00000001 0x00000000 | 0xF0 |
----------------------------------------------
| 0x8 | 0x00000000 0x00000001 | 0x0F |
| 0x8 | 0x00000001 0x00000000 | 0xF0 |
----------------------------------------------
| 0x9 | 0x00000000 0x000F4240 | 0x0F |
| 0x9 | 0x00000001 0x00000000 | 0xF0 |
----------------------------------------------
| 0xA | 0x00000000 0x00000001 | 0x0F |
| 0xA | 0x00000000 0x00000000 | 0xF0 |
----------------------------------------------
| 0xB | 0x00000000 0x00000000 | 0x0F |
| 0xB | 0x00000000 0x00000000 | 0xF0 |
----------------------------------------------
| 0xC | 0x00000000 0x20100000 | 0x0F |
| 0xC | 0x00000000 0x00000000 | 0xF0 |
----------------------------------------------
| 0xD | 0x00000000 0x20500000 | 0x0F |
| 0xD | 0x00000000 0x00000000 | 0xF0 |
----------------------------------------------
| 0xE | 0x00000000 0x20900000 | 0x0F |
| 0xE | 0x00000000 0x00000000 | 0xF0 |
----------------------------------------------
| 0x0 | 0x00000000 0x00000001 | 0x0F |
----------------------------------------------
Примечание:
byte_enable «выбирает» в какие байты регистра писать: так, в рамках самой первой транзакции записали 0x00000001 в нижние 32 бита регистра 0x5 (старшие 32 бита при этом не изменились).
Смотреть транзакции в SignalTap не всегда может быть удобно: на хосте есть возможность включить дополнительный дебаг через переменные окружения. Их можно подсмотреть в главе Troubleshooting Altera Stratix V Network Reference Platform Porting Guide.
Нам необходима переменная ACL_HAL_DEBUG. Выставляем её значение в 2 и запускаем хостовое приложение vector_add:
root@socfpga:~/myvectoradduint# export ACL_HAL_DEBUG=2
root@socfpga:~/myvectoradduint# ./vector_add
// <пропущен различный дебаг>
:: Launching kernel 0 on accelerator 0.
:: Writing inv image [ 0] @ 0x28 := 1
:: Writing inv image [ 4] @ 0x2c := f4240
:: Writing inv image [ 8] @ 0x30 := f4240
:: Writing inv image [12] @ 0x34 := 1
:: Writing inv image [16] @ 0x38 := 1
:: Writing inv image [20] @ 0x3c := 1
:: Writing inv image [24] @ 0x40 := 1
:: Writing inv image [28] @ 0x44 := 1
:: Writing inv image [32] @ 0x48 := f4240
:: Writing inv image [36] @ 0x4c := 1
:: Writing inv image [40] @ 0x50 := 1
:: Writing inv image [44] @ 0x54 := 0
:: Writing inv image [48] @ 0x58 := 0
:: Writing inv image [52] @ 0x5c := 0
:: Writing inv image [56] @ 0x60 := 20100000
:: Writing inv image [60] @ 0x64 := 0
:: Writing inv image [64] @ 0x68 := 20500000
:: Writing inv image [68] @ 0x6c := 0
:: Writing inv image [72] @ 0x70 := 20900000
:: Writing inv image [76] @ 0x74 := 0
:: Accelerator 0 reporting status 2.
:: Accelerator 0 is done.
Можно заметить, что 0x28 — это байтовый оффсет 5-го 64-битного регистра.
Видно, что адреса и данные совпадают, правда, в этом дебаге нет информации о транзакции в нулевой регистр (даже если ACL_HAL_DEBUG поставить равным пяти).
Результат настройки:
- work_dim — 0x1 — т.к. у нас одномерный вектор.
- workgroup_size — 0xF4240 или 1000000.
- global_size — 0xF4240 для первого измерения и 0x1 для всех остальных.
- num_groups — 0x1 для всех измерений.
- local_size — 0xF4240 для первого измерения и 0x1 для всех остальных.
- global_offset — 0x0 для всех измерений.
- input_x, input_y, input_z — 0x20100000, 0x20500000, 0x20900000 соответственно.
После настройки ядра дергается младший бит нулевого регистра, который производит старт вычислений.
Скрытый текст
Пока крутится барабан, хочу передать привет разработчикам из Альтеры.
Зачем вы включаете обработку через регистр с названием status?
Может стоило его назвать иначе?
Зачем вы включаете обработку через регистр с названием status?
Может стоило его назвать иначе?
Симуляция ядра
Теперь мы понимаем, как настраивается ядро — давайте же его просимулируем!
Очень удобно, что у ядра всего два интерфейса — один для настройки, другой для чтения данных (еще есть интерфейс (размером в один сигнал) для выставления прерывания — но нам это не так интересно).
Для симуляции ядра нам надо сделать всё как в реальной жизни:
- настроить его (теперь мы знаем что в какой последовательности надо писать, а интерфейс относительно простой).
- предоставить доступ до глобальной памяти, где выделены буфера и лежат данные.
Конечно же, никакой линукс и хостовое приложение симулировать мы не очень хотим, поэтому в качестве первого приближения можно ограничиться следующей схемой:
- cra_driver — драйвер для настройки ядра.
- vector_add_system — ядро, которое симулируем (DUT).
- avalon_mm_clock_crossing — переброс данных с частоты кернела (140 МГц) на частоту чтения из контроллера (100 МГц) и обратно.
- prepare_data — простые таски, которые записывают данные в буферы X и Y перед началом симуляции.
- avalon_mm_interconnect — мультиплексирование и арбитраж двух Avalon-MM интерфейсов.
- ddr3_contoller, ddr3_model — симуляционные модели Altera Hard Memory Controller и DDR3 памяти. Настройки модели и контроллера совпадают с теми, которые используются в модуле hps.
Драйвер для настройки ядра это просто последовательный вызов следующего таска с теми настройками ядра, что мы раскопали благодаря SignalTap'у:
task cra_write( input bit [3:0] _addr, bit [63:0] _data, bit [7:0] _byteenable );
$display("%m: _addr = 0x%x, _data = 0x%x, _byteenable = 0x%x",
_addr, _data, _byteenable );
@( posedge clk );
cra_addr <= _addr;
cra_wr_data <= _data;
cra_byteenable <= _byteenable;
cra_wr_en <= 1'b0;
@( posedge clk );
cra_wr_en <= 1'b1;
@( posedge clk );
cra_wr_en <= 1'b0;
// dummy waiting
repeat (10) @( posedge clk );
endtask
Скрытый текст
initial
begin
wait( ram_init_done );
wait( test_data_init_done );
cra_write( 4'h5, 64'h000F424000000000, 8'hF0 );
cra_write( 4'h5, 64'h0000000100000000, 8'hF0 );
cra_write( 4'h6, 64'h00000000000F4240, 8'h0F );
cra_write( 4'h6, 64'h0000000100000000, 8'hF0 );
cra_write( 4'h7, 64'h0000000000000001, 8'h0F );
cra_write( 4'h7, 64'h0000000100000000, 8'hF0 );
cra_write( 4'h8, 64'h0000000000000001, 8'h0F );
cra_write( 4'h8, 64'h0000000100000000, 8'hF0 );
cra_write( 4'h9, 64'h00000000000F4240, 8'h0F );
cra_write( 4'h9, 64'h0000000100000000, 8'hF0 );
cra_write( 4'hA, 64'h0000000000000001, 8'h0F );
cra_write( 4'hA, 64'h0000000000000000, 8'hF0 );
cra_write( 4'hB, 64'h0000000000000000, 8'h0F );
cra_write( 4'hB, 64'h0000000000000000, 8'hF0 );
cra_write( 4'hC, 64'h0000000020100000, 8'h0F );
cra_write( 4'hC, 64'h0000000000000000, 8'hF0 );
cra_write( 4'hD, 64'h0000000020500000, 8'h0F );
cra_write( 4'hD, 64'h0000000000000000, 8'hF0 );
cra_write( 4'hE, 64'h0000000020900000, 8'h0F );
cra_write( 4'hE, 64'h0000000000000000, 8'hF0 );
cra_write( 4'h0, 64'h0000000000000001, 8'h0F );
end
Для подготовки данных для которых необходимо сделать расчет пишем аналогичный таск, который будет писать в память по заранее забитым адресам. (Для выделения адресов нам никого (систему) спрашивать не надо). Для демонстрации не обязательно записывать 2x1000000 чисел как надо по условиям задачи — достаточно пару тысяч, для того, чтобы посмотреть как это работает. Если мы данные не запишем, то из памяти будут считываться x (unknown value), т.к. в эти ячейки ничего не было записано.
Смотрим результат симуляции (все числа на времянках в 16-ричной форме) (скриншоты лучше открыть в отдельном окне):
Приходит сигнал start и через несколько тактов оба LSU одновременно выставляют запрос на чтение данных с размером берста равным 0x10 = 16. Интересно, что сначала принимаются три запроса только от LSU_X, а потом от LSU_Y: это видно по сигналу avm_waitrequest запрос на чтение принимается только тогда, когда waitrequest равен нулю. Почему запросы не чередуются, как можно было бы ожидать — это вопрос к шедуллеру и интерконнекту до глобальной памяти.
Почему было именно три запроса?
Количество слов в FIFO в LSU равно 64, и можно было бы ожидать 4 запроса на чтение, но в там сделан такой трюк:
parameter READTHRESHOLD = FIFODEPTH - MAXBURSTCOUNT - 4;
assign too_many_reads_pending = (reads_pending + fifo_used) >= READTHRESHOLD; // make sure there are fewer reads posted than room in the FIFO
Мы отправили 3 запроса на чтение (3 x 16 = 48 слов): еще есть место для 16 слов, но чтение остановится пока не осводится четыре слова. Небольшого берста из 12 не будет, т.к. этот модуль так не умеет (да и смысла особого в этом нет — лишнюю логику тратить).
Через некоторое время приходят прочитанные данные (см. сигнал LSU_X_avm_readdatavalid) и почти сразу выставляется сигнал LSU_X_o_valid, извещающий, что 32-битные данные в LSU_X_o_readdata готовы для дальнейшей обработки, но нас затыкают: LSU_X_i_stall выставляется в единицу. Дело в том, что конвееру, который будет складывать данные не имеет данных от LSU_Y. Поэтому всё затыкается пока не придет LSU_Y_avm_readdatavalid и выставится LSU_Y_o_valid.
На следующий такт после этого приходит LSU_Z_i_valid, который извещает о вадидности LSU_Z_i_writedata: конвеер сложения отработал за один такт.
LSU_Z дожидается накопления необходимого количества данных: это конечно же 16 256-битных слов и производит транзакцию на запись. Параллельно с этим LSU_X и LSU_Y производят свои транзакции чтения по мере опустошения своих FIFO. Заметим, что LSU_X_i_stall и LSU_Y_i_stall не взводятся, а значит конвеер не затыкается и каждый такт принимает новые данные.
Посмотрев многое количество транзакции видно, что больше затыканий не происходит.
Вывод:
Из-за того, что конвеер никогда не останавливается (кроме самого начала) понятно, что он работает с максимальной производительностью, и именно он является узким звеном в этом простом примере. Так как чтение происходит по 256-битным словам, то есть смысл организовать параллельное сложение восьми 32-битных чисел — тогда скорее всего ботлнек произойдет в доступе к памяти. Большее количество сумматоров смысла особого не имеет
Благодаря симуляции мы можем вносить какие-то изменения к RTL-код (например, длину берста или величину кэшей) для того, чтобы посмотреть как это будет влиять на скорость вычислений и очень быстро увидеть ответ: для этого не надо будет ждать 10-15 минут для пересборки всего проекта Квартусом — достаточно запустить симуляцию в симуляторе.
Конечно, какие-то архитектурные изменения необходимо вносить через редактирование *.cl файла (например использовать специальные директивы): можно будет перегенерировать код, и так же запустить симуляцию и увидеть полученный выигрыш без сборки всего проекта.
Заключение
Мы узнали во что превращается ядро vector_add со стороны FPGA и как оно настраивается.
Как видим, многие вопросы, которые возникли бы у FPGA разработчиков (если бы писали с нуля), были решены:
- размеры кэшей
- архитектура
- сделан интерконнект и арбитраж интерфейсов
- автоматическая подстройка тактовой частоты под получившееся значение
Конечно, не факт, что то, что сделано автоматически — это самое оптимальное.
Скрытый текст
Хотя бы потому, что по дефолту настройки сборки Quartus'a (а именно она определяет всё, кроме архитектуры) были «сбалансированные».
В компиляторе есть огромное количество настроек, которые позволяют улучшить:
Конечно же, нельзя сразу соптимизировать все четыре компонента: хотите больше частоту — платите за это временем сборки и ресурсами FPGA.
В принципе, для того, чтобы выставить необходимые настройки необязательно быть опытным FPGA разработчиком, досточно следовать Advisor'ам, которые есть в GUI и просто накликать то, что он предлагает. Еще можно ознакомиться с видеоуроками (на сайте Altera).
В компиляторе есть огромное количество настроек, которые позволяют улучшить:
- время сборки
- количество ресурсов
- тактовую частоту
- энергопотребление
Конечно же, нельзя сразу соптимизировать все четыре компонента: хотите больше частоту — платите за это временем сборки и ресурсами FPGA.
В принципе, для того, чтобы выставить необходимые настройки необязательно быть опытным FPGA разработчиком, досточно следовать Advisor'ам, которые есть в GUI и просто накликать то, что он предлагает. Еще можно ознакомиться с видеоуроками (на сайте Altera).
Низкоуровневая оптимизация и ручной тюнинг — это самый последний шаг. Вы часто лезете в ассемблер после сборки gcc? Для начала надо провести высоуровневую оптимизацию в *.cl и поиграться с настройками Quartus'a.
Хорошо, что есть возможность отсимулировать ядро: можно прикинуть его производительность, не имея железа (платы) на руках.
Спасибо за внимание! Буду рад вопросам и замечаниям в комментариях или в личной почте.
Комментарии (6)
rPman
06.12.2015 18:00+1Я сломался где то по середине статьи, но если эта статься будет сохранена то это еще одна статья на разработку кода для FPGA на русском языке.
Пожалуйста, пишите такие статьи,… за этим будущее.
Надеюсь скоро тулчейн под это дело будет таким же простым, как сейчас собирается приложение под той же ms visual studio, вроде бы много доступно для изменений но по умолчанию никуда лезть не нужно.
Daffodil
13.12.2015 06:33Очень крутое исследование. Осталось понять что произойдет когда ресурсов FPGA не хватит для реализации kernel'a: он просто упадет с ошибкой, или попытается сгенерировать какой-нибудь сложный планировщик с выгрузкой промежуточных результатов в память.
ishevchuk
13.12.2015 11:36Спасибо за отзыв :)
Я не увидел сложного планировщика, когда ресурсов не хватило (я пробовал еще один дефолтный пример (перемножениее матриц)): в предварительном отчете перед запуском квартуса там фигурировали числа > 200%, и потом совершенно справедливо заругался Quartus. Возможно, когда числа очень близкие к 100% он проводит какие-то оптимизации, но я с этим не игрался.
Наверно, они могли бы это (сложный планировщик) сделать, но тогда пользователь получал бы маленькие числа в производительности и ругал бы технологию (хотя, вполне вероятно, что он просто плохо описал ядро). Возможно, сложные планировщики для уменьшения ресурсов не в их комерческих интересах: не влезает в этот чип — возьми жирнее/дороже (или трать своё время на ручную оптимизацию).
Daffodil
13.12.2015 06:46Насколько эта технология универсальна: можно ли выкачивать данные с других IP висящих на avalon шине?
Было бы удобно, если бы LSU накачивал данные непосредственно с I/O портов, к примеру с видеокамеры или сетевого порта, kernel'ы выполняли бы какую-нибудь предобработку (например сжатие), после чего другой LSU выгружал бы результаты в хостовую память.ishevchuk
13.12.2015 11:22+1Из-за того, что код сгенеренный получается открытым, особых проблем в подсоединении кернела в свою систему я не вижу, но надо поработать руками :)
Для перекачки данных с видеокамеры или сетевого порта Altera предлагает использовать I/O Channels.
Я глубо это не копал, но есть пример, где данные забираются с UDP-корки.
VBKesha
Спасибо за статью, теперь ясно как оно работает.