Altera + OpenCL: вскрываем ядро

Автор: admin от 6-12-2015, 17:10, посмотрело: 649

Altera + OpenCL: вскрываем ядро

Всем привет!

В прошлой статье я запустил простой 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-симуляторах?

  • Какие блоки занимают больше всего ресурсов? Можно ли как-то это соптимизировать?


Давайте взглянём на его внутренности! Добро пожаловать под кат!

Harnessing the Power of FPGAs using Altera’s OpenCL Compiler (осторожно, больше ста слайдов, ~16 МБ).

Altera + OpenCL: вскрываем ядро

Прошивка состоит из:

  • 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)


Altera + OpenCL: вскрываем ядро


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)


Altera + OpenCL: вскрываем ядро

Модули:

  • pll (0 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.

Примерная схема выглядит так:
Altera + OpenCL: вскрываем ядро

  • 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 получают «команды» на чтение данных и делают запросы к глобальной памяти.

  • Прочитанные данные буферизируются в памяти (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_STREAMINGLSU_X, LSU_Y (READ = 1, STYLE = «STREAMING»)

    • LSU_WRITE_STREAMINGLSU_Z (READ = 0, STYLE = «STREAMING»)



    LSU_READ_STREAMING

    Altera + OpenCL: вскрываем ядро
    Обратим внимание на параметры модуля:
    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

    Altera + OpenCL: вскрываем ядро
    Поступающие 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_bandwidthThe 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.



    К сожалению нет пояснений в каких единицах и как это считается: с одной стороны в профилировщике писалось 6400 MB/s, с другой стороны по расчетам 6400 MB/s никак не получаются: 400 (МГц, тактовая частота DDR) * 32 (бит, ширина сигнала данных на DDR-интерфейсе) * 2 (работа по двум фронтам) = 25600 Mb/s = 3200 MB/s. Либо надо считать в обе стороны?

    max_burstMaximum burst size for the slave interface.


    В нашем случае — 16, что дает BURSTCOUNT_WIDTH = 5. Но почему именно 16? Интерфейс fpga2hps_sdram поддерживает max_burstcount = 128. 16 — это какое-то магическое число, подходит всем? :)

    latencyAn 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.



    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 (через интерфейс реконфигурации). Как только реконфигурация закончилась на кернел уже подается нужная частота.



    Итого:

    • На ядро будет подана та, частота которую можно подавать. Если логика оказалась слишком ёмкой, и не удалось уложиться в заданное число, то вычисления не сломаются — просто они будут медленее идти.

    • Если же допустимая частота выше, то 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 после загрузки ядра и стартом вычислений, чтобы я успел подключиться с помощью дебаггера).

    Altera + OpenCL: вскрываем ядро


    Как управляется ядро


    Интерфейс 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:
    [email protected]:~/myvectoradduint# export ACL_HAL_DEBUG=2
    [email protected]:~/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 соответственно.


    После настройки ядра дергается младший бит нулевого регистра, который производит старт вычислений.



    Симуляция ядра


    Теперь мы понимаем, как настраивается ядро — давайте же его просимулируем!
    Очень удобно, что у ядра всего два интерфейса — один для настройки, другой для чтения данных (еще есть интерфейс (размером в один сигнал) для выставления прерывания — но нам это не так интересно).

    Для симуляции ядра нам надо сделать всё как в реальной жизни:

    • настроить его (теперь мы знаем что в какой последовательности надо писать, а интерфейс относительно простой).

    • предоставить доступ до глобальной памяти, где выделены буфера и лежат данные.


    Конечно же, никакой линукс и хостовое приложение симулировать мы не очень хотим, поэтому в качестве первого приближения можно ограничиться следующей схемой:
    Altera + OpenCL: вскрываем ядро

    • cra_driver — драйвер для настройки ядра.

    • vector_add_system — ядро, которое симулируем (Altera + OpenCL: вскрываем ядро


      Приходит сигнал 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: конвеер сложения отработал за один такт.
      Altera + OpenCL: вскрываем ядро


      LSU_Z дожидается накопления необходимого количества данных: это конечно же 16 256-битных слов и производит транзакцию на запись. Параллельно с этим LSU_X и LSU_Y производят свои транзакции чтения по мере опустошения своих FIFO. Заметим, что LSU_X_i_stall и LSU_Y_i_stall не взводятся, а значит конвеер не затыкается и каждый такт принимает новые данные.

      Altera + OpenCL: вскрываем ядро


      Посмотрев многое количество транзакции видно, что больше затыканий не происходит.

      Вывод:
      Из-за того, что конвеер никогда не останавливается (кроме самого начала) понятно, что он работает с максимальной производительностью, и именно он является узким звеном в этом простом примере. Так как чтение происходит по 256-битным словам, то есть смысл организовать параллельное сложение восьми 32-битных чисел — тогда скорее всего ботлнек произойдет в доступе к памяти. Большее количество сумматоров смысла особого не имеет

      Благодаря симуляции мы можем вносить какие-то изменения к RTL-код (например, длину берста или величину кэшей) для того, чтобы посмотреть как это будет влиять на скорость вычислений и очень быстро увидеть ответ: для этого не надо будет ждать 10-15 минут для пересборки всего проекта Квартусом — достаточно запустить симуляцию в симуляторе.

      Конечно, какие-то архитектурные изменения необходимо вносить через редактирование *.cl файла (например использовать специальные директивы): можно будет перегенерировать код, и так же запустить симуляцию и увидеть полученный выигрыш без сборки всего проекта.

      Заключение


      Мы узнали во что превращается ядро vector_add со стороны FPGA и как оно настраивается.

      Как видим, многие вопросы, которые возникли бы у FPGA разработчиков (если бы писали с нуля), были решены:

      • размеры кэшей

      • архитектура

      • сделан интерконнект и арбитраж интерфейсов

      • автоматическая подстройка тактовой частоты под получившееся значение


      Конечно, не факт, что то, что сделано автоматически — это самое оптимальное.


      Низкоуровневая оптимизация и ручной тюнинг — это самый последний шаг. Вы часто лезете в ассемблер после сборки gcc? Для начала надо провести высоуровневую оптимизацию в *.cl и поиграться с настройками Quartus'a.

      Хорошо, что есть возможность отсимулировать ядро: можно прикинуть его производительность, не имея железа (платы) на руках.

      Спасибо за внимание! Буду рад вопросам и замечаниям в комментариях или в личной почте.

      Источник: Хабрахабр

Категория: Программирование

Уважаемый посетитель, Вы зашли на сайт как незарегистрированный пользователь.
Мы рекомендуем Вам зарегистрироваться либо войти на сайт под своим именем.

Добавление комментария

Имя:*
E-Mail:
Комментарий:
Полужирный Наклонный текст Подчеркнутый текст Зачеркнутый текст | Выравнивание по левому краю По центру Выравнивание по правому краю | Вставка смайликов Выбор цвета | Скрытый текст Вставка цитаты Преобразовать выбранный текст из транслитерации в кириллицу Вставка спойлера
Введите два слова, показанных на изображении: *