- 武蔵野 Advent Calendar 18日目 の記事です。
概要
OpenCL for FPGA
- ホスト PC に FPGA が接続されているシステムにおいて、両者の挙動（下記２種類のコード）を一括して OpenCL C 言語で記述できる環境がある。
- OpenCL ホストコード
- C コンパイラに通され、ホスト CPU で動作する。
- OpenCL カーネルコード
- 高位合成をサポートしたツールチェーンに通され、ディジタル回路設計に変換し FPGA 上に書き込まれる。
- OpenCL 規格は Khronos Group
によって制定されているものであり、これに従っている限り、かつ動作速度を気にしなければ、GPU 用のコードを各社 FPGA でも使用できるはず
- 多くの場合高速ないし高電力効率な演算が目的とされており動作速度は重要なのだが、まず"動く"ところから始められる事は開発のハードルを下げてくれる
Intel FPGA SDK for OpenCL
- Intel 社（旧 Altera 社） FPGA 用の OpenCL for FPGA 環境
- 当 SDK はこれまで有償製品だったが、v17.1 アップデート（2017/11/06）よりライセンス不要となった
-
Intel FPGA SDK for OpenCL Version 17.1 Release Notes によれば :
Removed requirement for IntelR FPGA SDK for OpenCL license. You can also now run your OpenCL kernel without a paid runtime licenseとのこと。 - なお、Intel 社 FPGA 向けの FPGA コンフィグ作成は、エントリ向け FPGA（Cyclone シリーズ等）は無償、ハイエンド FPGA（Stratix シリーズ等）は有償ライセンスとなっている。
- つまり、エントリ向け FPGA を搭載したボードと対応 BSP があれば、OpenCL for FPGA 開発が無償で始められるようになった。（下表）
-
Intel FPGA SDK for OpenCL Version 17.1 Release Notes
| エントリ向け FPGA | ハイエンド FPGA | |
|---|---|---|
| FPGA コンフィグ作成 | 無償 | 有償 |
| OpenCL 開発環境 |
|
|
- 動作には FPGA ボードに対応した BSP (Board Support Package) が必要
- BSP : FPGA ボード上の外部 I/O 等、周辺回路を OpenCL カーネルから共通の規格で叩けるようにする、「FPGA 上の回路として動かすデバイスドライバ」のようなもの。
- ボードメーカーなどが供給する前提だが、自作することも可能。
-
Terasic 社 DE10-Standard 開発キット などはエントリ向け FPGA 搭載品でありボードメーカーによる BSP が提供されている
- しかし、搭載 FPGA は ARM CPU との混載（SoC）タイプであり、ハード・ソフト協調動作と言ってもワンチップ上で動く形になる。
- PC との組み合わせ由来の相性問題などを気にせず使用できる点では SoC FPGA を利用する利点が大きい
- 【筆者の趣味】普通の PC に装着していかにも「外付け」感たっぷりに使いたい。かつ趣味の範囲（ライセンス無料）でやりたい。
Cyclone V GT FPGA 開発キット（FPGAボード）
- エントリ向け FPGA (Cyclone V) を搭載しておりライセンス無償で開発可能
- PCI Express x4 (gen2) + DDR3 SDRAM 装備
- OpenCL 対応要件が揃っている。そのため非公式ながら有志により BSP が作成されている
-
"電波望遠鏡用の分光器をAltera SDK for OpenCL使ってサクッと作ってみた" (slideshare) などに OpenCL での使用例あり
- OpenCL 対応要件が揃っている。そのため非公式ながら有志により BSP
- イーサネットなどの外部 I/O を備える
- 非公式 BSP では外部 I/O 操作は未対応
- 【筆者の興味】なので外部 I/O 操作を当キットの BSP に追加できると面白い。
host pipe 機能
- Intel FPGA SDK for OpenCL v17.1 から host pipe 機能 (デバイス側 DRAM へのデータストアを介せず、直接データをやり取りする機能)がサポートされた
- 通常 OpenCL においてはホスト⇔デバイス間のデータやり取りをデバイス側に装備された DRAM 経由で行うが、リアルタイム性向上などの理由で直接 DMA で転送できるようにする拡張機能
- 当機能は BSP 側に予約された外部 I/O ポート（実際はホストとの DMA やりとりを行うポートにつながっている）を呼び出す形で実装されており、利用には BSP 側の対応が必要と考えられる ( 参考 : Altera Wiki の記事 )
- FPGA ボード側の RAM 量に限度があるローコスト・低消費電力向けの用途においても生きるかもしれない
- 【筆者の興味】これも Cyclone V GT FPGA 開発キットの BSP に追加できると面白い。
本記事の取り組み
- Cyclone V GT FPGA 開発キットを使い始める（Lチカ編）
- Cyclone V GT FPGA 開発キットを使い始める（イーサネット編）
- 非公式 OpenCL BSP を動かしてみる
- OpenCL BSP に LED 操作ポートをつけて直接操作する
- （未完）OpenCL BSP にイーサネットポートをつけて操作する
1. Cyclone V GT FPGA 開発キットを使い始める（Lチカ編）
- 目標 : ボード単独でLチカする
-
Intel Quartus Prime Lite 17.1 (Quartus-lite-17.1.0.590-linux.tar) をダウンロードしてきてインストール
- 以下では
/opt/intelFPGA_liteにインストールしたものとする -
$ quartusしてメイン画面が開く事を確認
- 以下では
-
Debian 9 (Stretch) で動かすときのワークアラウンド
- デバイス命名則を ethX 形式に戻す :
/etc/default/grubにGRUB_CMDLINE_LINUX="net.ifnames=0"を追加して$ sudo update-grub - jessie を
/etc/apt/sources.listに追加（stretch より下の行に書くか、apt-pinning して優先順位を下げる）して、# apt-get update、# apt-get install libpng12-0 - multi-arch の設定
# dpkg --add-architecture i386、# apt-get updateのあと、# apt-get install libc6-i386 libx11-6:i386 libxext6:i386 libxft2:i386 libncurse5:i386 libpng12-0:i386 -
# dpkg-reconfigure localesから en_US.UTF-8 を追加生成ロケールに指定（デフォルトロケールは好きなものでよい） - 公式サポートに Ubuntu 16.04.2 LTS があるのでこれを使えば苦労が無いと思われる
- デバイス命名則を ethX 形式に戻す :
-
PATH を通す用の参考スクリプト
setup_intelFPGA_lite_17_1.sh
# usage : source {this_file.sh} {path_to_install_dir (ex. /opt/intelFPGA/17.1)}
#LICENSE_FILE=/opt/intelFPGA/license/{license_file}.dat
#DEFAULT_BSP=c5gt_eth_custom
if [ -z "$1" ]; then
echo "usage : source {this_file.sh} {path_to_install_dir (ex. /opt/intelFPGA/17.1)}"
else
export QUARTUS_ROOTDIR="$1"/quartus
export MODELSIM_ASE_DIR="$1"/modelsim_ase
export MODELSIM_AE_DIR="$1"/modelsim_ae
export INTELFPGAOCLSDKROOT="$1"/hld
export ALTERAOCLSDKROOT="$INTELFPGAOCLSDKROOT"
export SOPC_KIT_NIOS2="$QUARTUS_ROOTDIR"/../nios2eds
export PATH=$PATH:"$QUARTUS_ROOTDIR"/bin:"$QUARTUS_ROOTDIR"/linux64:"$QUARTUS_ROOTDIR"/../qsys/bin:"$QUARTUS_ROOTDIR"/sopc_builder/bin:"$QUARTUS_ROOTDIR"/sopc_builder/model/bin:"$SOPC_KIT_NIOS2"/bin:"$INTELFPGAOCLSDKROOT"/linux64/bin:"$INTELFPGAOCLSDKROOT"/bin:"$MODELSIM_AE_DIR"/bin:"$MODELSIM_ASE_DIR"/bin
export AOCL_BOARD_PACKAGE_ROOT="$INTELFPGAOCLSDKROOT"/board/"$DEFAULT_BSP"
export LD_LIBRARY_PATH="$AOCL_BOARD_PACKAGE_ROOT"/linux64/lib:"$INTELFPGAOCLSDKROOT"/host/linux64/lib
export QUARTUS_64BIT=1
#if [ -z "$LM_LICENSE_FILE" ]; then
# export LM_LICENSE_FILE="$LICENSE_FILE"
#else
# export LM_LICENSE_FILE=$LM_LICENSE_FILE:"$LICENSE_FILE"
#fi
fi
- USB ケーブルドライバ導入
-
/etc/udev/rules.d/51-usbblaster.rulesに以下記入して再起動（もしくは udevadm を使用） -
$ quartus_pgmwして [Hardware Setup] の欄に USB-Blaster II が出ていたら OK
-
/etc/udev/rules.d/51-usbblaster.rules
# Altera USB-Blaster for Quartus FPGA Software
ACTION=="add|change", SUBSYSTEM=="usb", ATTRS{idVendor}=="09fb", ATTRS{idProduct}=="6010|6810", MODE="0666"
-
サポートページ から Kit Installation をダウンロードして展開
-
cycloneVGT_5cgtfd9ef35_fpga_v13.0.0.1/examples/golden_top/c5gt_pro_goldentop.vのdefine USERのコメントアウトを解除（モジュール引数最後のコンマは処理）し、以下を追記
c5gt_pro_goldentop.v（追記部分）
reg [24:0] cnt_r;
reg led_r;
assign user_led[0] = led_r;
assign user_led[1] = led_r;
assign user_led[2] = led_r;
assign user_led[3] = led_r;
assign user_led[4] = led_r;
assign user_led[5] = led_r;
assign user_led[6] = led_r;
assign user_led[7] = led_r;
always @ (posedge clkin_50 or negedge cpu_resetn) begin
if (cpu_resetn == 1'b0) begin
cnt_r <= 25'd24999999;
led_r <= 1'b0;
end else if (cnt_r == 25'd0) begin
cnt_r <= 25'd24999999;
led_r <= ~led_r;
end else begin
cnt_r <= cnt_r - 25'd1;
led_r <= led_r;
end
end
- quartus で
c5gt_pro_goldentop.qpfを開き、[Task] -> [Assembler (Generate programming files)]-
output_files/c5gt_pro_goldentop.sofが生成される
-
- コンフィグレーション
- [Tools] -> [Programmer] を起動し、Hardware setup から認識されている USB-Blaster II を選択
- [Auto Detect] の後、厳密なデバイス名を選択するように指示されるので 5CGTFD9E5 選択
- 5CGTFD9E5 デバイスの上で右クリックして [Change File]、
output_files/c5gt_pro_goldentop.sofを選択 - [Program / Configure] チェックボックスをオン
- [Start] をクリック
- HSMC ポート A の近くにある 4個 x 2列 並んだ LED が１秒に１回点滅することを確認
2. Cyclone V GT FPGA 開発キットを使ってみる（イーサネット編）
- 目標 : イーサネットポートを開通させる
- 初期状態の golden_top プロジェクトを準備
IP コア準備
- イーサネット関連の IP コアを Quartus の IP Catalog から準備してプロジェクトに追加
-
cycloneVGT_5cgtfd9ef35_fpga_v13.0.0.1/examples/board_update_portalで使われているイーサネット関係の IP コアと同じものを選び、IP コアのコンフィグもそれに倣う- 但し [alignment packet to 32bit boundary] と [full duplex flow control] は外した
- Triple-Speed Ethernet -> eth_1g_mac の名前で生成
- ALTDDIO_OUT -> eth_1g_clkbuf の名前で生成
- Altera PLL -> eth_1g_pll の名前で生成
自分で記述する部分
- 最低限の MAC コンフィグを行うコントローラをプロジェクトに追加 (
cycloneVGT_5cgtfd9ef35_fpga_v13.0.0.1/examples/golden_top/eth_1g_mac_init.v)
eth_1g_mac_init.v
//------------------------------------------------------------------------------
// eth_1g_mac_init.v
//------------------------------------------------------------------------------
module eth_1g_mac_init(
input clk,
input xrst,
output [7:0] reg_addr,
input [31:0] reg_data_read,
output [31:0] reg_data_write,
output reg_rval,
output reg_wval,
input reg_busy
);
// 命令セット
// 4bit 8bit 32bit
// {inst, addr, immidiate}
parameter INST_WAIT = 4'd0;
parameter INST_WRITE = 4'd1;
parameter WAIT_INFTY = 32'hFFFF_FFFF;
wire [43:0] curr_inst;
wire [3:0] curr_inst_type;
wire [7:0] curr_inst_addr;
wire [31:0] curr_inst_immediate;
wire end_curr_inst;
reg [7:0] prog_count_r;
reg start_curr_inst_r;
reg [31:0] wait_count_r;
reg [7:0] reg_addr_r;
reg [31:0] reg_data_write_r;
reg reg_wval_r;
// 命令列メモリ
function [43:0] inst_mem_f;
input [7:0] prog_count_arg;
case (prog_count_r)
// リセット後数サイクル待つ
8'd0 : inst_mem_f = { INST_WAIT, 8'h00, 32'd10};
// COMMAND CONFIG レジスタ
//8'd1 : inst_mem_f = {INST_WRITE, 8'h02, 32'h0B00_0033}; // 10Mbps
//8'd1 : inst_mem_f = {INST_WRITE, 8'h02, 32'h0900_0033}; // 100Mbps
8'd1 : inst_mem_f = {INST_WRITE, 8'h02, 32'h0900_003B}; // GbE
// 停止
8'd2 : inst_mem_f = { INST_WAIT, 8'h00, WAIT_INFTY};
default : inst_mem_f = { INST_WAIT, 8'h00, WAIT_INFTY};
endcase
endfunction
// 今見ている命令
assign curr_inst = inst_mem_f(prog_count_r);
assign curr_inst_type = curr_inst[43:40];
assign curr_inst_addr = curr_inst[39:32];
assign curr_inst_immediate = curr_inst[31:0];
// 現命令が当サイクルで終了し、次のサイクルで次の命令へ進む
function end_curr_inst_f;
input [3:0] curr_inst_type_arg;
input reg_busy_arg;
input [31:0] wait_count_arg;
case (curr_inst_type_arg)
INST_WAIT : end_curr_inst_f = ((|wait_count_arg) == 1'b0);
INST_WRITE : end_curr_inst_f = ~reg_busy_arg;
default : end_curr_inst_f = 1'b0;
endcase
endfunction
assign end_curr_inst = end_curr_inst_f(curr_inst_type,
reg_busy,
wait_count_r);
// プログラムカウンタ
always @ (posedge clk or negedge xrst) begin
if (xrst == 1'b0) begin
prog_count_r <= 8'd0;
end else if (end_curr_inst == 1'b1) begin
// 現命令の終了条件に合致していればプログラムカウンタを増やす
prog_count_r <= prog_count_r + 8'd1;
end else begin
prog_count_r <= prog_count_r;
end
end
// 命令開始時に1サイクルだけ上がるスタートフラグ
always @ (posedge clk or negedge xrst) begin
if (xrst == 1'b0) begin
start_curr_inst_r <= 1'b1;
end else if (start_curr_inst_r == 1'b0 && end_curr_inst == 1'b1) begin
start_curr_inst_r <= 1'b1;
end else begin
start_curr_inst_r <= 1'b0;
end
end
// スタートフラグが立ったらアドレスと書き込み用即値を読み込む
always @ (posedge clk or negedge xrst) begin
if (xrst == 1'b0) begin
reg_addr_r <= 8'd0;
reg_data_write_r <= 32'd0;
end else if (start_curr_inst_r == 1'b1) begin
reg_addr_r <= curr_inst_addr;
reg_data_write_r <= curr_inst_immediate;
end else begin
reg_addr_r <= reg_addr_r;
reg_data_write_r <= reg_data_write_r;
end
end
// 書き込みフラグ動作
always @ (posedge clk or negedge xrst) begin
if (xrst == 1'b0) begin
reg_wval_r <= 1'b0;
end else if (end_curr_inst) begin
reg_wval_r <= 1'b0;
end else if (start_curr_inst_r == 1'b1 && curr_inst_type == INST_WRITE) begin
reg_wval_r <= 1'b1;
end else begin
reg_wval_r <= reg_wval_r;
end
end
// 残り WAIT サイクル数を示すカウンタ
always @ (posedge clk or negedge xrst) begin
if (xrst == 1'b0) begin
wait_count_r <= WAIT_INFTY;
end else if (start_curr_inst_r == 1'b1 && curr_inst_type == INST_WAIT) begin
wait_count_r <= curr_inst_immediate;
end else if ((&wait_count_r) == 1'b1) begin
// wait_count_r が WAIT_INFTY の時はデクリメントしない
wait_count_r <= wait_count_r;
end else begin
wait_count_r <= wait_count_r - 32'd1;
end
end
// 出力ポート
assign reg_addr = reg_addr_r;
assign reg_data_write = reg_data_write_r;
assign reg_rval = 1'b0;
assign reg_wval = reg_wval_r;
endmodule
-
cycloneVGT_5cgtfd9ef35_fpga_v13.0.0.1/examples/golden_top/c5gt_pro_goldentop.vのdefine ETHERNET、define USERのコメントアウトを解除（モジュール引数最後のコンマは処理）し、以下を追記
c5gt_pro_goldentop.v（追記部分）
// MDIO
wire enet_mdio_oen;
wire enet_mdio_out;
// PHY clock
wire eth_clk_125;
wire eth_clk_25;
wire eth_clk_2p5;
wire eth_clk_locked;
wire eth_clk_selected;
// 1G MAC control signal
wire eth_mode;
wire ena_10;
wire [7:0] mac_reg_addr;
wire [31:0] mac_reg_data_read;
wire [31:0] mac_reg_data_write;
wire mac_reg_rval;
wire mac_reg_wval;
wire mac_reg_busy;
// 1G MAC internal loopback
wire [31:0] l4_data;
wire l4_startofpacket;
wire l4_endofpacket;
wire [1:0] l4_empty;
wire l4_ready;
wire l4_valid;
assign enet_mdio = ( ! enet_mdio_oen ) ? enet_mdio_out : 1'bz;
assign eth_clk_selected = (eth_mode) ? eth_clk_125 : // GbE Mode = 125MHz clock
(ena_10) ? eth_clk_2p5 : // 10Mb Mode = 2.5MHz clock
eth_clk_25; // 100Mb Mode = 25MHz clock
eth_1g_clkbuf eth_1g_clkbuf_inst (
.aclr ( ! cpu_resetn),
.datain_h (1'b1),
.datain_l (1'b0),
.outclock (eth_clk_selected),
.dataout (enet_gtx_clk)
);
eth_1g_pll eth_1g_pll_inst(
.refclk (clkin_50), // input refclk.clk
.rst (~cpu_resetn), // input reset.reset
.outclk_0 (eth_clk_125), // output outclk0.clk
.outclk_1 (eth_clk_25), // output outclk1.clk
.outclk_2 (eth_clk_2p5), // output outclk2.clk
.locked (eth_clk_locked) // output locked.export
);
eth_1g_mac_init mac_init_inst(
.clk (clkin_50),
.xrst (cpu_resetn),
.reg_addr (mac_reg_addr), // output [7:0]
.reg_data_read (mac_reg_data_read), // input [31:0]
.reg_data_write(mac_reg_data_write), // output [31:0]
.reg_rval (mac_reg_rval), // output
.reg_wval (mac_reg_wval), // output
.reg_busy (mac_reg_busy) // input
);
eth_1g_mac eth_1g_mac_inst(
.clk (clkin_50), // input control_port_clock_connection.clk
.reset (~cpu_resetn), // input reset_connection.reset
.reg_addr (mac_reg_addr), // input [7:0] control_port.address
.reg_data_out (mac_reg_data_read), // output [31:0] .readdata
.reg_rd (mac_reg_rval), // input .read
.reg_data_in (mac_reg_data_write), // input [31:0] .writedata
.reg_wr (mac_reg_wval), // input .write
.reg_busy (mac_reg_busy), // output .waitrequest
.tx_clk (eth_clk_selected), // input pcs_mac_tx_clock_connection.clk
.rx_clk (enet_rx_clk), // input pcs_mac_rx_clock_connection.clk
.set_10 (1'b0), // input mac_status_connection.set_10 (ignored)
.set_1000 (1'b0), // input .set_1000 (ignored)
.eth_mode (eth_mode), // output .eth_mode
.ena_10 (ena_10), // output .ena_10
.rgmii_in (enet_rx_d), // input [3:0] mac_rgmii_connection.rgmii_in
.rgmii_out (enet_tx_d), // output [3:0] .rgmii_out
.rx_control (enet_rx_dv), // input .rx_control
.tx_control (enet_tx_en), // output .tx_control
.ff_rx_clk (clkin_50), // input receive_clock_connection.clk
.ff_tx_clk (clkin_50), // input transmit_clock_connection.clk
.ff_rx_data (l4_data), // output [31:0] receive.data
.ff_rx_sop (l4_startofpacket), // output .startofpacket
.ff_rx_eop (l4_endofpacket), // output .endofpacket
.ff_rx_mod (l4_empty), // output [1:0] .empty
.ff_rx_rdy (l4_ready), // input .ready
.ff_rx_dval (l4_valid), // output .valid
.rx_err (), // output [5:0] .error
.ff_tx_data (l4_data), // input [31:0] transmit.data
.ff_tx_sop (l4_startofpacket), // input .startofpacket
.ff_tx_eop (l4_endofpacket), // input .endofpacket
.ff_tx_mod (l4_empty), // input [1:0] .empty
.ff_tx_rdy (l4_ready), // output .ready
.ff_tx_wren (l4_valid), // input .valid
.ff_tx_err (1'b0), // input .error
.mdc (enet_mdc), // output mac_mdio_connection.mdc
.mdio_in (enet_mdio), // input .mdio_in
.mdio_out (enet_mdio_out), // output .mdio_out
.mdio_oen (enet_mdio_oen), // output .mdio_oen
.ff_tx_crc_fwd (1'b0), // input mac_misc_connection.ff_tx_crc_fwd
.ff_tx_septy (), // output .ff_tx_septy
.tx_ff_uflow (), // output .tx_ff_uflow
.ff_tx_a_full (), // output .ff_tx_a_full
.ff_tx_a_empty (), // output .ff_tx_a_empty
.rx_err_stat (), // output [17:0] .rx_err_stat
.rx_frm_type (), // output [3:0] .rx_frm_type
.ff_rx_dsav (), // output .ff_rx_dsav
.ff_rx_a_full (), // output .ff_rx_a_full
.ff_rx_a_empty () // output .ff_rx_a_empty
);
//----------------------------------------------------------------------
// indicate eth_mode
assign user_led[6] = ~eth_mode;
assign user_led[7] = ~ena_10;
//----------------------------------------------------------------------
// PHY Power On Reset
reg [19:0] epcount_r; // PHY interface: need minimum 10ms delay for POR
always @(posedge clkin_50 or negedge cpu_resetn) begin
if (cpu_resetn == 1'b0) begin
epcount_r <= 20'd0;
end else if (epcount_r[19] == 1'b0) begin
epcount_r <= epcount_r + 20'd1;
end else begin
epcount_r <= epcount_r;
end
end
assign enet_resetn = !epcount_r[18];
//----------------------------------------------------------------------
// packet count to LED
reg [5:0] pktcount_r;
assign user_led[5:0] = ~pktcount_r;
always @ (posedge clkin_50 or negedge cpu_resetn) begin
if (cpu_resetn == 1'b0) begin
pktcount_r <= 6'd0;
end else if (l4_endofpacket & l4_valid) begin
pktcount_r <= pktcount_r + 6'd1;
end else begin
pktcount_r <= pktcount_r;
end
end
コンパイルと動作検証
- 固定 IP を振った適当な PC と 1 対 1 接続し、コンパイル -> コンフィグを行う
- Ethernet PHY の auto negotiation の結果を見ない簡易構成なので、リンクアップした速度モードは PHY チップの横に 3 個速度モードを示す LED があるのでそれを読み取る。
- その後 init モジュールの中で MAC 側に手動で指定、必要であれば再コンパイル -> コンフィグ。
- 何かパケットを投げると、パケット1個毎に USER LED のうち下位 6 ビットがカウンタに割り当ててあるので受信パケットの数だけインクリメントされていく
3. 非公式 OpenCL BSP を動かしてみる
- 目標 : PCIe を通してボードから Hello World をさせる
ツールキットセットアップ
- Intel Quartus Prime Standard 17.1 (Quartus-17.1.0.590-linux-complete.tar 使用) を
/opt/intelFPGAにインストール（OpenCL SDK のインストールをオンにすること）- OpenCL SDK は Standard 版しかないので、
/opt/intelFPGA_lite/17.1以下で# ln -s /opt/intelFPGA/17.1/hld ./hldとしてリンクを張る
- OpenCL SDK は Standard 版しかないので、
-
The board support package of Cyclone V GT Development Kit for Intel FPGA SDK OpenCL からダウンロード、展開
- /opt/intelFPGA/17.1/hld/board/ にも c5gt ディレクトリごとコピーしておく
- BSP 配置テスト :
$ aoc -list-boardsで c5gt が出てくれば OK
PCIe ドライバコンパイルと導入 (debian 9 (stretch) の場合)
-
/usr/src/linux-headers-xxxx-amd64を/usr/src/kernels/{`uname -r`}にソフトリンク -
c5gt/linux64/driver以下のaclpci_cmd.cを編集
aclpci_cmd.c
// 230 行目付近
// 誤：
ret = get_user_pages(target_task, target_task->mm,
start_page + got * PAGE_SIZE,
num_pages - got, 1, 1,
p + got, vma);
// 正：
ret = get_user_pages_remote(target_task, target_task->mm,
start_page + got * PAGE_SIZE,
num_pages - got, 1,
p + got, vma);
-
モジュールコンパイル
-
# aocl installは CentOS 用のモジュールインストール手順を試みるため、Debian 上で実行してエラー終了した場合ファイルを削除してしまう。 -
# aocl installの代わりにc5gt/linux64/driver以下で$ ./make_all.sh -
aclpci_drv.koが生成されれば OK
-
-
モジュールインストール
-
/lib/modules/xxxx-amd64/kernel/misc/aclpci_drv.koにコピーし、$ sudo depmod -a
-
-
ユーザからのアクセス権付与（
/etc/udev/rules.d/52-aclpci_drv.rulesに以下記入）
/etc/udev/rules.d/52-aclpci_drv.rules
KERNEL=="acl0", MODE="0666"
-
BSP を flash 抜きでコンパイル
- readme.txt にあるように base.flash を flash に焼いてもいいが、BSP 改造で数をこなしたいのでこの手順を試す
- c5gt/hardware/c5gt の top.qpf を開いて、Platform Designer から system.qsys を開いて [Generate HDL], その後 quartus 本体で [Assembler (Generate Programming Files)] 実行
- top.sof が生成される
-
PCIe ドライバの動作テスト
- top.sof を焼き込み、再起動（warm reboot）
- PC のマザーボードによっては再起動時に PCIe のリセットが行われない場合もある模様。その場合は AC アダプタで電源を維持したままシャットダウン、リブート
- 再起動後、
$ aocl diagnoseで DIAGNOSTIC_PASSED と出れば成功
Hello World
-
Hello World サンプル をダウンロード、展開
- カーネル作成
$ aoc device/hello_world.cl -o bin/hello_world.aocx -board=c5gt- bin/hello_world/top.sof を quartus_pgmw で焼き込み、warm reboot
- ホスト動作
-
$ make: /bin/host が生成される - partial reconfiguration による CvP が使えないので、
CL_CONTEXT_COMPILER_MODE_INTELFPGA=3 ./hostとして CvP を飛ばして起動 （参考ページ : マクニカ社サイトの FAQ）
-
- 結果
$ CL_CONTEXT_COMPILER_MODE_INTELFPGA=3 ./host
Querying platform for info:
==========================
CL_PLATFORM_NAME = Intel(R) FPGA SDK for OpenCL(TM)
CL_PLATFORM_VENDOR = Intel(R) Corporation
CL_PLATFORM_VERSION = OpenCL 1.0 Intel(R) FPGA SDK for OpenCL(TM), Version 17.1
Querying device for info:
========================
CL_DEVICE_NAME = c5gt : CycloneV GT FPGA Development Kit
CL_DEVICE_VENDOR = Intel(R) Corporation
CL_DEVICE_VENDOR_ID = 4466
CL_DEVICE_VERSION = OpenCL 1.0 Intel(R) FPGA SDK for OpenCL(TM), Version 17.1
CL_DRIVER_VERSION = 17.1
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 = 402653184
CL_DEVICE_IMAGE_SUPPORT = true
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 = 100663296
CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS = 3
CL_DEVICE_MEM_BASE_ADDR_ALIGN = 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.
- FPGA 側の回路としてハードコーディングされた
Hello from Altera's OpenCL Compiler!の文字列を PCIe 経由で読み出して表示するサンプルであり、最低限の動作は確認できた
4. OpenCL BSP に LED 操作ポートをつけて直接操作する
- 目標 : OpenCL から再度Lチカをする
-
FPGAマガジン No.14 第8章 "画像データをダイレクト入力！ OpenCLによるカメラ制御&画像処理" にて DE0-nano-SoC 搭載の Cyclone V SoC 上における例が示されており、全体の参考になる
BSP への channel 追加
- c5gt/hardware/c5gt/board_spec.xml に追加 channel を定義
c5gt/hardware/c5gt/board_spec.xml
<channels>
<interface name="board" port="user_led_ctrl_kernel_in" type="streamsink" width="8" chan_id="ch_user_led_ctrl"/>
</channels>
- c5gt/hardware/c5gt/user_led_ctrl/user_led_ctrl.v
- channel 出力（avalon-ST）を受けて LED の状態 8 bit の信号線状態を制御するモジュール
c5gt/hardware/c5gt/user_led_ctrl/user_led_ctrl.v
module user_led_ctrl (
input kernel_clk,
input kernel_xrst,
input [7:0] kernel_in_data,
output kernel_in_ready,
input kernel_in_valid,
output [7:0] led_out
);
reg [7:0] led_out_r;
assign kernel_in_ready = 1'b1;
assign led_out = led_out_r;
always @ (posedge kernel_clk or negedge kernel_xrst) begin
if (kernel_xrst == 1'b0) begin
led_out_r <= 8'd0;
end else if (kernel_in_valid) begin
led_out_r <= kernel_in_data;
end else begin
led_out_r <= led_out_r;
end
end
endmodule
- user_led_ctrl.v を platform designer (旧 qsys) から扱えるようにするラッパー作成
- platform designer の IP Catalog で [new..] を押して作成、name / display_name を user_led_ctrl とし、user_led_ctrl.v を [files] のタブに追加
- [Signals & Interfaces] で以下のようなポート割り当てを実施して保存
- board.qsys への結線
- platform designer から board.qsys を開き、user_led_ctrl を追加
- kernel_clk, kernel_reset を board.qsys 内で配線
- led_out, kernel_in を export
- 以下配線例
-
system.qsys への配線
- board.qsys を変更したので board モジュールのポートが 2 個増える
- user_led_ctrl_kernel_in はそのまま空けておく (OpenCL カーネルコンパイル後に自動結線される)
- user_led_ctrl_led_out は同名で export
-
c5gt/hardware/c5gt/user_led_ctrl/top.v 変更
- 外部 IF の user_led が [3:0] になっているので [7:0] に変更（ピンアサインは qsf に予約済み）
- system モジュールに user_led_ctrl_led_out ポートが追加されるので、現在
assign user_led = 4'b0101;となっているところをコメントアウトして、代わりに user_led[7:0] を user_led_ctrl_led_out ポートに結線
プログラム
- LED テスト用カーネルコード
led_test.cl
#pragma OPENCL EXTENSION cl_intel_channels : enable
#include "led_test.h"
channel uchar ch_user_led_ctrl __attribute__((depth(0))) __attribute__((io("ch_user_led_ctrl")));
__attribute__((max_global_work_dim(0)))
__kernel void user_led_ctrl (__global ushort * restrict led_out) {
write_channel_intel(ch_user_led_ctrl, (*led_out));
}
- LED テスト用ホストコード例
led_test.cpp
#include <assert.h>
#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <cstring>
#include "AOCLUtils/aocl_utils.h"
#include <sys/types.h>
#include <sys/time.h>
#include <sys/socket.h>
#include <netinet/in.h>
#include <arpa/inet.h>
#include <netdb.h>
#include <sched.h>
#include <unistd.h>
using namespace aocl_utils;
typedef unsigned char uchar;
// OpenCL runtime configuration
#define STRING_BUFFER_LEN 1024
#define KERNEL_AOCX_NAME "kernel_led_test"
// カーネル定義、カーネル制御コマンド群をロード
#define __OPENCL_HOST__
#include "led_test.h"
// OpenCL globals
static cl_platform_id platform = NULL;
static cl_device_id device = NULL;
static cl_context context = NULL;
static cl_program program = NULL;
static cl_kernel kernels[K_NUM_KERNELS];
static cl_command_queue queues[K_NUM_KERNELS];
// User globals
cl_mem cl_user_led_ctrl_led_out;
// Function prototypes
void check_clCreateBuffer (cl_mem* memptr, cl_mem_flags flags, size_t size);
void check_clSetKernelArg (ushort kernel_id, cl_uint arg_index, cl_mem* memptr);
void check_clFinish (ushort kernel_id);
bool init();
void cleanup();
static void device_info_ulong ( cl_device_id device, cl_device_info param, const char* name);
static void device_info_uint ( cl_device_id device, cl_device_info param, const char* name);
static void device_info_bool ( cl_device_id device, cl_device_info param, const char* name);
static void device_info_string ( cl_device_id device, cl_device_info param, const char* name);
static void display_device_info( cl_device_id device );
int main() {
// OpenCL カーネル初期化
if(!init()) {
return -1;
}
// OpenCL カーネル引数用のメモリ領域確保
check_clCreateBuffer(&cl_user_led_ctrl_led_out, CL_MEM_READ_ONLY, sizeof(uchar));
// 確保した領域をカーネル引数にセット
check_clSetKernelArg(K_USER_LED_CTRL, 0, &cl_user_led_ctrl_led_out);
int led_out = 0x00;
for (int i = 0; i < 20; i ++) {
cl_int status;
status = clEnqueueWriteBuffer(queues[K_USER_LED_CTRL], cl_user_led_ctrl_led_out, CL_TRUE, 0, sizeof(uchar), &led_out, 0, NULL, NULL);
checkError(status, "Failed to write cmd to %s\n", kernel_names[K_USER_LED_CTRL]);
status = clEnqueueTask(queues[K_USER_LED_CTRL], kernels[K_USER_LED_CTRL], 0, NULL, NULL);
checkError(status, "Failed to enqueue %s\n", kernel_names[K_USER_LED_CTRL]);
check_clFinish(K_USER_LED_CTRL);
led_out = (~led_out);
usleep(500000);
}
// OpenCL カーネル解放
cleanup();
return 0;
}
// チェック付きの clCreateBuffer
void check_clCreateBuffer(cl_mem* memptr, cl_mem_flags flags, size_t size) {
cl_int status;
(*memptr) = clCreateBuffer(context, flags, size, NULL, &status);
checkError(status, "Failed to create buffer");
return;
}
// チェック付きの clSetKernelArg
void check_clSetKernelArg(ushort kernel_id, cl_uint arg_index, cl_mem* memptr) {
cl_int status;
status = clSetKernelArg(kernels[kernel_id], arg_index, sizeof(cl_mem), (void*)memptr);
checkError(status, "Failed to set %s kernel argument %u\n", kernel_names[kernel_id], arg_index);
return;
}
// チェック付きの clFinish
void check_clFinish(ushort kernel_id) {
cl_int status;
status = clFinish(queues[kernel_id]);
checkError(status, "Failed to finish %s\n", kernel_names[kernel_id]);
return;
}
/////// OPENCL HELPER FUNCTIONS ///////
bool init() {
cl_int status;
// Start everything at NULL to help identify errors
for(int i = 0; i < K_NUM_KERNELS; ++i){
kernels[i] = NULL;
queues[i] = NULL;
}
// Locate files via. relative paths
if(!setCwdToExeDir()) {
return false;
}
// Get the OpenCL platform.
platform = findPlatform("Intel(R) FPGA");
if(platform == NULL) {
printf("ERROR: Unable to find Intel(R) FPGA OpenCL platform.\n");
return false;
}
// User-visible output - Platform information
//{
// char char_buffer[STRING_BUFFER_LEN];
// printf("Querying platform for info:\n");
// printf("==========================\n");
// clGetPlatformInfo(platform, CL_PLATFORM_NAME, STRING_BUFFER_LEN, char_buffer, NULL);
// printf("%-40s = %s\n", "CL_PLATFORM_NAME", char_buffer);
// clGetPlatformInfo(platform, CL_PLATFORM_VENDOR, STRING_BUFFER_LEN, char_buffer, NULL);
// printf("%-40s = %s\n", "CL_PLATFORM_VENDOR ", char_buffer);
// clGetPlatformInfo(platform, CL_PLATFORM_VERSION, STRING_BUFFER_LEN, char_buffer, NULL);
// printf("%-40s = %s\n\n", "CL_PLATFORM_VERSION ", char_buffer);
//}
// Query the available OpenCL devices.
scoped_array<cl_device_id> devices;
cl_uint num_devices;
devices.reset(getDevices(platform, CL_DEVICE_TYPE_ALL, &num_devices));
// We'll just use the first device.
device = devices[0];
// Display some device information.
display_device_info(device);
// Create the context.
context = clCreateContext(NULL, 1, &device, &oclContextCallback, NULL, &status);
checkError(status, "Failed to create context");
// Create the command queues
for(int i = 0; i < K_NUM_KERNELS; i ++) {
queues[i] = clCreateCommandQueue(context, device, CL_QUEUE_PROFILING_ENABLE, &status);
checkError(status, "Failed to create command queue (%d)", i);
}
// Create the program.
std::string binary_file = getBoardBinaryFile(KERNEL_AOCX_NAME, device);
printf("Using AOCX: %s\n", binary_file.c_str());
program = createProgramFromBinary(context, binary_file.c_str(), &device, 1);
// Build the program that was just created.
status = clBuildProgram(program, 0, NULL, "", NULL, NULL);
checkError(status, "Failed to build program");
// Create the kernel - name passed in here must match kernel name in the
// original CL file, that was compiled into an AOCX file using the AOC tool
for(int i = 0; i < K_NUM_KERNELS; i ++) {
kernels[i] = clCreateKernel(program, kernel_names[i], &status);
checkError(status, "Failed to create kernel (%d: %s)", i, kernel_names[i]);
}
return true;
}
// Free the resources allocated during initialization
void cleanup() {
for(int i = 0; i < K_NUM_KERNELS; i ++) {
if (kernels[i]) {
clReleaseKernel(kernels[i]);
}
}
if (program) {
clReleaseProgram(program);
}
for(int i = 0; i<K_NUM_KERNELS; i ++) {
if (queues[i]) {
clReleaseCommandQueue(queues[i]);
}
}
if (context) {
clReleaseContext(context);
}
}
// Helper functions to display parameters returned by OpenCL queries
static void device_info_ulong( cl_device_id device, cl_device_info param, const char* name) {
cl_ulong a;
clGetDeviceInfo(device, param, sizeof(cl_ulong), &a, NULL);
printf("%-40s = %lu\n", name, a);
}
static void device_info_uint( cl_device_id device, cl_device_info param, const char* name) {
cl_uint a;
clGetDeviceInfo(device, param, sizeof(cl_uint), &a, NULL);
printf("%-40s = %u\n", name, a);
}
static void device_info_bool( cl_device_id device, cl_device_info param, const char* name) {
cl_bool a;
clGetDeviceInfo(device, param, sizeof(cl_bool), &a, NULL);
printf("%-40s = %s\n", name, (a?"true":"false"));
}
static void device_info_string( cl_device_id device, cl_device_info param, const char* name) {
char a[STRING_BUFFER_LEN];
clGetDeviceInfo(device, param, STRING_BUFFER_LEN, &a, NULL);
printf("%-40s = %s\n", name, a);
}
// Query and display OpenCL information on device and runtime environment
static void display_device_info( cl_device_id device ) {
printf("Querying device for info:\n");
printf("========================\n");
device_info_string(device, CL_DEVICE_NAME, "CL_DEVICE_NAME");
device_info_string(device, CL_DEVICE_VENDOR, "CL_DEVICE_VENDOR");
device_info_uint(device, CL_DEVICE_VENDOR_ID, "CL_DEVICE_VENDOR_ID");
device_info_string(device, CL_DEVICE_VERSION, "CL_DEVICE_VERSION");
device_info_string(device, CL_DRIVER_VERSION, "CL_DRIVER_VERSION");
device_info_uint(device, CL_DEVICE_ADDRESS_BITS, "CL_DEVICE_ADDRESS_BITS");
device_info_bool(device, CL_DEVICE_AVAILABLE, "CL_DEVICE_AVAILABLE");
device_info_bool(device, CL_DEVICE_ENDIAN_LITTLE, "CL_DEVICE_ENDIAN_LITTLE");
device_info_ulong(device, CL_DEVICE_GLOBAL_MEM_CACHE_SIZE, "CL_DEVICE_GLOBAL_MEM_CACHE_SIZE");
device_info_ulong(device, CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE, "CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE");
device_info_ulong(device, CL_DEVICE_GLOBAL_MEM_SIZE, "CL_DEVICE_GLOBAL_MEM_SIZE");
device_info_bool(device, CL_DEVICE_IMAGE_SUPPORT, "CL_DEVICE_IMAGE_SUPPORT");
device_info_ulong(device, CL_DEVICE_LOCAL_MEM_SIZE, "CL_DEVICE_LOCAL_MEM_SIZE");
device_info_ulong(device, CL_DEVICE_MAX_CLOCK_FREQUENCY, "CL_DEVICE_MAX_CLOCK_FREQUENCY");
device_info_ulong(device, CL_DEVICE_MAX_COMPUTE_UNITS, "CL_DEVICE_MAX_COMPUTE_UNITS");
device_info_ulong(device, CL_DEVICE_MAX_CONSTANT_ARGS, "CL_DEVICE_MAX_CONSTANT_ARGS");
device_info_ulong(device, CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE, "CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE");
device_info_uint(device, CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, "CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS");
device_info_uint(device, CL_DEVICE_MEM_BASE_ADDR_ALIGN, "CL_DEVICE_MEM_BASE_ADDR_ALIGN");
device_info_uint(device, CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE, "CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE");
device_info_uint(device, CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR, "CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR");
device_info_uint(device, CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT, "CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT");
device_info_uint(device, CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT, "CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT");
device_info_uint(device, CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG, "CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG");
device_info_uint(device, CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT, "CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT");
device_info_uint(device, CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE, "CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE");
{
cl_command_queue_properties ccp;
clGetDeviceInfo(device, CL_DEVICE_QUEUE_PROPERTIES, sizeof(cl_command_queue_properties), &ccp, NULL);
printf("%-40s = %s\n", "Command queue out of order? ", ((ccp & CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE)?"true":"false"));
printf("%-40s = %s\n", "Command queue profiling enabled? ", ((ccp & CL_QUEUE_PROFILING_ENABLE)?"true":"false"));
}
}
- 両者から読まれるインクルードファイル
led_test.h
#ifndef __LED_TEST_H__
#define __LED_TEST_H__
// 両者で必要になるカーネル ID
enum KERNELS {
#ifdef __SIM__
#endif
K_USER_LED_CTRL,
K_NUM_KERNELS
};
#ifdef __OPENCL_HOST__
// ホストで必要になるカーネル ID とカーネル名の対応リスト
static const char* kernel_names[K_NUM_KERNELS] =
{
#ifdef __SIM__
#endif
"user_led_ctrl"
};
#else
#endif
#endif
動作検証
-
コンパイル、動作
- あらかじめ hello_world サンプルの common を同ディレクトリに配置しておく
$ aoc -g -v -profile -I/opt/intelFPGA_lite/17.1/hld/include/kernel_headers -o kernel_led_test.aocx -board=c5gt -report ../kernel/led_test.cl$ g++ -fPIC -I./common/inc `aocl compile-config` `aocl link-config` -o led_test ./common/src/AOCLUtils/*.cpp ./led_test.cppCL_CONTEXT_COMPILER_MODE_INTELFPGA=3 ./led_test
-
結果動画
Cyclone V GT Dev Kit の OpenCL BSP から LED を操作できるようにした pic.twitter.com/0JpmiMH0GY
— homelith (@homelith) 2017年12月17日
- ホストから 8 bit uchar の値をデバイス側に送信し、その 8 bit を LED 状態に反映するように記述された OpenCL カーネルをキックしている
- これを 500 ミリ秒ごとに値変更することで点滅させた
5. （未完）OpenCL BSP にイーサネットポートをつけて操作する
-
目標 : Ethernet ポートを OpenCL から叩けるようになる
-
問題発生
が、Intel社 の Triple-speed Ethernet を試用モード（OpenCore Plus Evaluation Mode）で使用する場合、time_limited.sof のみ生成でき、rbf や aocx を生成できない。
（最初は flash への書き込みができなくても、元々 CvP 無しで quartus_pgmw で焼き込む前提のフローなら動作可能かも、と思って取り組み始めたのだが、途中でホストプログラムが aocx を要求することに気付いた）
何か試用版のままで動作させるワークアラウンドがあるかも知れないが今のところ思いついていない .. -
今後どうするか
OpenCores 10_100_1000 Mbps tri-mode ethernet MAC （PHY インタフェースは GMII）と Quartus に搭載されている GMII to RGMII コンバータを使用すればライセンスの制約を受けずに開通させられそうに思う、が詳しくは未調査
まとめ
- Intel FPGA SDK for OpenCL v17.1 がライセンス無償化されたのを機に、PCI Express 接続された Cyclone V GT FPGA 開発キットの外部直接 I/O を OpenCL でプログラミングする第一歩を実施できた
- "5. OpenCL BSP にイーサネットポートをつけて操作する" については、今後 OpenCores 10_100_1000 Mbps tri-mode ethernet MAC を使用できないか取り組み、アップデートしていきたい
- host channel 対応可能かどうかも今後検討したい
- おうち SmartNIC への道は遠い