• Пару слов об FPGA
• На чём запускать?
• Процесс разработки (workflow)
• OpenCL BSP
• Компилируем кернел
• … и запускаем
• Заключение

Пару слов об FPGA (ПЛИС)

На чём запускать?

PCIe

SoC

Процесс разработки (workflow)

• Код кернела описывается в файле *.cl.
• Готовится хостовое приложение на С/C++, которое будет производить выделение необходимых объемов памяти и «загрузку» значений в кернел.
• С помощью утилиты aoc, которая входит в Altera OpenCL SDK, «компилируется» ядро в aocx файл. С помощью gcc собирается хостовое приложение.
• При запуске host_app произойдет загрузка прошивки FPGA, в ядро загрузятся подготовленные данные и начнется их обработка.
• Счетчики для профилирования собирают данные, которые поместятся в файл profile.mon.
• С помощью утилиты aocl можно посмотреть этот отчет и сделать вывод: удовлетворяет ли по времени выполнения/производительности эта реализация.
• Если удовлетворяет, то можно перекомпилировать ядро без --profile, т.к. профилирующие счетчики отнимают ресурсы в FPGA. С другой стороны, если дополнительных ядер не планируется добавлять, то можно и не пересобирать.
• Если не удовлетворяет, то надо оптимизировать/писать ручками/взять другой чип или смириться.

Сборка aocx

• kernel.cl скармливается clang, который переводит описание в IR, а так же проводит различные оптимизации.
• Генерируется RTL-ное Verilog IP-ядро. Сгенеренные файлы доступны для чтения (незашифрованы) и могут быть просимулировать в обычном симуляторе (например, ModelSim). Однако, там не весь код автосгенеренный: есть модули, которые явно писал человек.
• Полученное IP «присоединяется» к дефолтому проекту для платы и получается обычный проект для Quartus'a.
• Проходит сборка проекта (Analysis & Synthesis, Fitter, Assembler). Именно этот пункт занимает наибольшее время (от десяти минут до нескольких часов): поиск оптимальных мест расположения примитивов требует много вычислений.
• Результат сборки, информация о борде и прочее размещают в aocx, который является просто ELF-файлом.

DE1-SoC OpenCL BSP

Скрытый текст

• Образ флешки (с неё загрузится linux).
• Дефолтный проект, где уже настроены все пины, а так же интерфейсы (fpga2sdram, lwhps2fpga и др.).
• Простенькие примеры.

root@socfpga:~# uname -a
Linux socfpga 3.13.0-00298-g3c7cbb9-dirty #3 SMP Fri Jul 4 15:42:32 CST 2014 armv7l GNU/Linux

root@socfpga:~# cat /etc/issue     
Poky 8.0 (Yocto Project 1.3 Reference Distro) 1.3 \n \l

root@socfpga:~# cat /proc/cpuinfo
processor       : 0
model name      : ARMv7 Processor rev 0 (v7l)
Features        : swp half thumb fastmult vfp edsp thumbee neon vfpv3 tls vfpd32
CPU implementer : 0x41
CPU architecture: 7
CPU variant     : 0x3
CPU part        : 0xc09
CPU revision    : 0

processor       : 1
model name      : ARMv7 Processor rev 0 (v7l)
Features        : swp half thumb fastmult vfp edsp thumbee neon vfpv3 tls vfpd32
CPU implementer : 0x41
CPU architecture: 7
CPU variant     : 0x3
CPU part        : 0xc09
CPU revision    : 0

Hardware        : Altera SOCFPGA
Revision        : 0000
Serial          : 0000000000000000

Run "source ./init_opencl.sh" to setup OpenCL Run-Time Environment, including loading driver, on this board. 
Do it once right after booting the board.

OpenCL Run-Time Environment is pre-installed in opencl_arm32_rte folder.

root@socfpga:~# cat init_opencl.sh
export ALTERAOCLSDKROOT=/home/root/opencl_arm32_rte
export AOCL_BOARD_PACKAGE_ROOT=$ALTERAOCLSDKROOT/board/c5soc
export PATH=$ALTERAOCLSDKROOT/bin:$PATH
export LD_LIBRARY_PATH=$ALTERAOCLSDKROOT/host/arm32/lib:$LD_LIBRARY_PATH
insmod $AOCL_BOARD_PACKAGE_ROOT/driver/aclsoc_drv.ko

root@socfpga:~/helloworld# ./helloworld
Querying platform for info:
==========================
CL_PLATFORM_NAME                         = Altera SDK for OpenCL
CL_PLATFORM_VENDOR                       = Altera Corporation
CL_PLATFORM_VERSION                      = OpenCL 1.0 Altera SDK for OpenCL, Version 14.0

Querying device for info:
========================
CL_DEVICE_NAME                           = de1soc_sharedonly : Cyclone V SoC Development Kit
CL_DEVICE_VENDOR                         = Altera Corporation
CL_DEVICE_VENDOR_ID                      = 4466
CL_DEVICE_VERSION                        = OpenCL 1.0 Altera SDK for OpenCL, Version 14.0
CL_DRIVER_VERSION                        = 14.0
CL_DEVICE_ADDRESS_BITS                   = 64
CL_DEVICE_AVAILABLE                      = true
CL_DEVICE_ENDIAN_LITTLE                  = true
CL_DEVICE_GLOBAL_MEM_CACHE_SIZE          = 32768
CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE      = 0
CL_DEVICE_GLOBAL_MEM_SIZE                = 536870912
CL_DEVICE_IMAGE_SUPPORT                  = false
CL_DEVICE_LOCAL_MEM_SIZE                 = 16384
CL_DEVICE_MAX_CLOCK_FREQUENCY            = 1000
CL_DEVICE_MAX_COMPUTE_UNITS              = 1
CL_DEVICE_MAX_CONSTANT_ARGS              = 8
CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE       = 134217728
CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS       = 3
CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS       = 8192
CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE       = 1024
CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR    = 4
CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT   = 2
CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT     = 1
CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG    = 1
CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT   = 1
CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE  = 0
Command queue out of order?              = false
Command queue profiling enabled?         = true
Using AOCX: hello_world.aocx

Kernel initialization is complete.
Launching the kernel...

Thread #2: Hello from Altera's OpenCL Compiler!

Kernel execution is complete.

Установка SDK

• Quartus
• Altera SDK for OpenCL
• SoC Embedded Design Suite: набор тулзов для сборки и отладки приложений на SoC'e.

• Могут потребоваться рутовые права, причем об этом вам скажут только в конце установки.
• После установки необходимо кое-чего прописывать в PATH, ALTERAOCLSDKROOT, QUARTUS_ROOTDIR. Что туда прописывать можно подчерпнуть из соответствующих гайдов.

export PATH=/home/ish/altera/14.1/quartus/bin:$PATH
export PATH=/home/ish/altera/14.1/hld/bin:$PATH
export PATH=/usr/local/DS-5/bin:$PATH
export PATH=/usr/local/DS-5/sw/gcc/bin:$PATH
export PATH=/home/ish/altera/14.1/hld/linux64/bin/:$PATH
export ALTERAOCLSDKROOT=/home/ish/altera/14.1/hld/
export QUARTUS_ROOTDIR=/home/ish/altera/14.1/quartus/
export LD_LIBRARY_PATH=/home/ish/altera/14.1/hld/linux64/lib/:$LD_LIBRARY_PATH

# необходимость в этой строчке появится чуть позже, но я указал эту переменную вместе с остальными
export AOCL_BOARD_PACKAGE_ROOT=/home/ish/altera/14.1/hld/board/de1soc

note:before the below operations,make sure you have install the opencl SDK 14.0 and SoCEDS 14.0.
1. directly unzip the de1soc_openCL_bsp.zip into %ALTERAOCLSDKROOT%/board directory.
2. set the "User variables" AOCL_BOARD_PACKAGE_ROOT to %ALTERAOCLSDKROOT%/board/de1soc
3. open the windows command window and type "aoc --list-boards", it should output "de1soc_sharedonly"

ish@xmr:~$ aoc --list-boards
Board list:
  de1soc_sharedonly

Собираем пример

// 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];
}

• пытается распознать, какие есть OpenCL девайсы
• перепрограммирует FPGA, используя aocx-файл
• инициализирует буфера для массивов X, Y, Z
• генерирует данные в массивах X и Y, а так же вычисляет (на процессоре) референсный ответ
• передает указатели на массивы в кернел
• запускает обработку
• дожидается её окончания
• сравнивает референсный ответ с тем, что посчитал кернел

ish@xmr:~/tmp/cl/vector_add$ aoc device/vector_add.cl -o bin/vector_add.aocx --board de1soc_sharedonly --profile -v
aoc: Environment checks are completed successfully.
You are now compiling the full flow!!
aoc: Selected target board de1soc_sharedonly
aoc: Running OpenCL parser....
aoc: OpenCL parser completed successfully.
aoc: Compiling....
aoc: Linking with IP library ...
aoc: First stage compilation completed successfully.
aoc: Hardware generation completed successfully.

+-------------------------------------------------------------------------------+
; Fitter Summary                                                                ;
+---------------------------------+---------------------------------------------+
; Fitter Status                   ; Successful - Sat Oct 17 21:36:01 2015       ;
; Quartus II 64-Bit Version       ; 14.1.0 Build 186 12/03/2014 SJ Full Version ;
; Revision Name                   ; top                                         ;
; Top-level Entity Name           ; top                                         ;
; Family                          ; Cyclone V                                   ;
; Device                          ; 5CSEMA5F31C6                                ;
; Timing Models                   ; Final                                       ;
; Logic utilization (in ALMs)     ; 5,570 / 32,070 ( 17 % )                     ;
; Total registers                 ; 9685                                        ;
; Total pins                      ; 103 / 457 ( 23 % )                          ;
; Total virtual pins              ; 0                                           ;
; Total block memory bits         ; 127,344 / 4,065,280 ( 3 % )                 ;
; Total DSP Blocks                ; 0 / 87 ( 0 % )                              ;
; Total HSSI RX PCSs              ; 0                                           ;
; Total HSSI PMA RX Deserializers ; 0                                           ;
; Total HSSI TX PCSs              ; 0                                           ;
; Total HSSI PMA TX Serializers   ; 0                                           ;
; Total PLLs                      ; 2 / 6 ( 33 % )                              ;
; Total DLLs                      ; 1 / 4 ( 25 % )                              ;
+---------------------------------+---------------------------------------------+

ish@xmr:~/tmp/cl/vector_add$ readelf -a bin/vector_add.aocx 
ELF Header:
  Magic:   7f 45 4c 46 01 01 01 00 00 00 00 00 00 00 00 00 
  Class:                             ELF32
  Data:                              2's complement, little endian
  Version:                           1 (current)
  OS/ABI:                            UNIX - System V
  ABI Version:                       0
  Type:                              NONE (None)
  Machine:                           Advanced Micro Devices X86-64
  Version:                           0x1
  Entry point address:               0x0
  Start of program headers:          0 (bytes into file)
  Start of section headers:          2370388 (bytes into file)
  Flags:                             0x0
  Size of this header:               52 (bytes)
  Size of program headers:           0 (bytes)
  Number of program headers:         0
  Size of section headers:           40 (bytes)
  Number of section headers:         20
  Section header string table index: 1

Section Headers:
  [Nr] Name              Type            Addr     Off    Size   ES Flg Lk Inf Al
  [ 0]                   NULL            00000000 000000 000000 00      0   0  0
  [ 1] .shstrtab         STRTAB          00000000 000080 00011c 00   S  0   0 128
  [ 2]                   PROGBITS        00000000 000200 001000 00      0   0 128
  [ 3] .acl.board        PROGBITS        00000000 001200 000011 00      0   0 128
  [ 4] .acl.compileoptio PROGBITS        00000000 001280 000002 00      0   0 128
  [ 5] .acl.version      PROGBITS        00000000 001300 00000a 00      0   0 128
  [ 6] .acl.file.0       PROGBITS        00000000 001380 000030 00      0   0 128
  [ 7] .acl.source.0     PROGBITS        00000000 001400 0006c2 00      0   0 128
  [ 8] .acl.nfiles       PROGBITS        00000000 001b00 000001 00      0   0 128
  [ 9] .acl.source       PROGBITS        00000000 001b80 0006c2 00      0   0 128
  [10] .acl.opt.rpt.xml  PROGBITS        00000000 002280 000019 00      0   0 128
  [11] .acl.mav.json     PROGBITS        00000000 002300 00107f 00      0   0 128
  [12] .acl.area.json    PROGBITS        00000000 003380 0009da 00      0   0 128
  [13] .acl.profiler.xml PROGBITS        00000000 003d80 002f08 00      0   0 128
  [14] .acl.profile_base PROGBITS        00000000 006d00 0009c8 00      0   0 128
  [15] .acl.autodiscover PROGBITS        00000000 007700 000071 00      0   0 128
  [16] .acl.autodiscover PROGBITS        00000000 007780 00021e 00      0   0 128
  [17] .acl.board_spec.x PROGBITS        00000000 007a00 0003eb 00      0   0 128
  [18] .acl.fpga.bin     PROGBITS        00000000 007e00 23ab98 00      0   0 128
  [19] .acl.quartus_repo PROGBITS        00000000 242a00 000151 00      0   0 128
Key to Flags:
  W (write), A (alloc), X (execute), M (merge), S (strings), l (large)
  I (info), L (link order), G (group), T (TLS), E (exclude), x (unknown)
  O (extra OS processing required) o (OS specific), p (processor specific)

There are no section groups in this file.

There are no program headers in this file.

There are no relocations in this file.

There are no unwind sections in this file.

No version information found in this file.

Запускаем кернел

root@socfpga:~/myvectoradduint# ls -l
-rwxr-xr-x    1 root     root         42525 Apr 16 06:57 vector_add
-rw-r--r--    1 root     root       2371188 Apr 16 06:58 vector_add.aocx

root@socfpga:~/myvectoradduint# ./vector_add 
Initializing OpenCL
Platform: Altera SDK for OpenCL
Using 1 device(s)
  de1soc_sharedonly : Cyclone V SoC Development Kit
Using AOCX: vector_add.aocx
Launching for device 0 (1000000 elements)

Time: 112.475 ms
Kernel time (device 0): 7.270 ms

Verification: PASS

root@socfpga:~/myvectoradduint# ls -l
-rw-r--r--    1 root     root           170 Apr 16 06:58 profile.mon
-rwxr-xr-x    1 root     root         42525 Apr 16 06:57 vector_add
-rw-r--r--    1 root     root       2371188 Apr 16 06:58 vector_add.aocx

ish@xmr:~/tmp/cl/vector_add$ aocl report bin/vector_add.aocx profile.mon

ish@xmr:~/tmp/cl/vector_add$ aocl vis bin/vector_add.aocx 

Заключение

• Не позволяет описывать процессы с точностью до такта (но ведь от этого мы и хотели уйти!)
• Неприменим на маленьких чипах: инфраструктура отъедает огромное количество ресурсов.

• Исходники примера и автосгеренеренный Verilog-код
• Altera SDK for OpenCL: Overview
• Altera SDK for OpenCL: Getting Started Guide
• Harnessing the Power of FPGAs using Altera’s OpenCL Compiler: очень большая презентация (>100 слайдов) для тех кто любит картинки смотреть :)
• Altera SDK for OpenCL: Programming Guide
• Altera SDK for OpenCL: Optimization Guide
• Altera SDK for OpenCL: Best Practices Guide