Cyclone V GT FPGA 開発キットの外部 I/O を OpenCL から叩く


- [武蔵野 Advent Calendar](https://qiita.com/advent-calendar/2017/musashino) 18日目 の記事です。

![P_20171217_214052_2.jpg](https://qiita-image-store.s3.amazonaws.com/0/216085/ae449020-d7da-891a-81dd-5e06eb5c9bd9.jpeg)

# 概要

#### OpenCL for FPGA
- ホスト PC に FPGA が接続されているシステムにおいて、両者の挙動（下記２種類のコード）を一括して OpenCL C 言語で記述できる環境がある。
- OpenCL ホストコード
    + C コンパイラに通され、ホスト CPU で動作する。
- OpenCL カーネルコード
    + 高位合成をサポートしたツールチェーンに通され、ディジタル回路設計に変換し FPGA 上に書き込まれる。
- OpenCL 規格は [Khronos Group :arrow_upper_right:](https://www.khronos.org/) によって制定されているものであり、これに従っている限り、かつ動作速度を気にしなければ、GPU 用のコードを各社 FPGA でも使用できるはず
    + 多くの場合高速ないし高電力効率な演算が目的とされており動作速度は重要なのだが、まず"動く"ところから始められる事は開発のハードルを下げてくれる

#### [Intel FPGA SDK for OpenCL :arrow_upper_right:](https://www.altera.co.jp/products/design-software/embedded-software-developers/opencl/overview.html)
- Intel 社（旧 Altera 社） FPGA 用の OpenCL for FPGA 環境
- 当 SDK はこれまで有償製品だったが、v17.1 アップデート（2017/11/06）よりライセンス不要となった
    + [Intel FPGA SDK for OpenCL Version 17.1 Release Notes :arrow_upper_right:](https://www.altera.co.jp/documentation/ewa1412772636144.html) によれば :  `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 開発が無償で始められるようになった。（下表）

|                     |  エントリ向け FPGA |    ハイエンド FPGA |
|:-------------------:|:------------------:|:------------------:|
| FPGA コンフィグ作成 |               無償 |               有償 |
|     OpenCL 開発環境 | ~~有償~~ ⇒ _無償_ | ~~有償~~ ⇒ _無償_ |


- 動作には FPGA ボードに対応した BSP (Board Support Package) が必要
    + BSP : FPGA ボード上の外部 I/O 等、周辺回路を OpenCL カーネルから共通の規格で叩けるようにする、「FPGA 上の回路として動かすデバイスドライバ」のようなもの。
    + ボードメーカーなどが供給する前提だが、自作することも可能。
- [Terasic 社 DE10-Standard 開発キット :arrow_upper_right:](www.terasic.com.tw/cgi-bin/page/archive.pl?Language=English&CategoryNo=205&No=1081&PartNo=1) などはエントリ向け FPGA 搭載品でありボードメーカーによる BSP が提供されている
    + しかし、搭載 FPGA は ARM CPU との混載（SoC）タイプであり、ハード・ソフト協調動作と言ってもワンチップ上で動く形になる。
    + PC との組み合わせ由来の相性問題などを気にせず使用できる点では SoC FPGA を利用する利点が大きい
    + 【筆者の趣味】普通の PC に装着していかにも「外付け」感たっぷりに使いたい。かつ趣味の範囲（ライセンス無料）でやりたい。

#### [Cyclone V GT FPGA 開発キット（FPGAボード） :arrow_upper_right:](https://www.altera.co.jp/products/boards_and_kits/dev-kits/altera/kit-cyclone-v-gt.html)
- エントリ向け FPGA (Cyclone V) を搭載しておりライセンス無償で開発可能
- PCI Express x4 (gen2) + DDR3 SDRAM 装備
    + OpenCL 対応要件が揃っている。そのため非公式ながら有志により [BSP :arrow_upper_right:](http://www.alterawiki.com/wiki/The_board_support_package_of_Cyclone_V_GT_Development_Kit_for_Intel_FPGA_SDK_OpenCL) が作成されている
    + ["電波望遠鏡用の分光器をAltera SDK for OpenCL使ってサクッと作ってみた" (slideshare) :arrow_upper_right:](https://www.slideshare.net/HirokiNakahara1/ss-63032205) などに OpenCL での使用例あり
- イーサネットなどの外部 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 の記事 :arrow_upper_right:](http://www.alteraforum.com/forum/showthread.php?t=57113) )
- FPGA ボード側の RAM 量に限度があるローコスト・低消費電力向けの用途においても生きるかもしれない
- 【筆者の興味】これも Cyclone V GT FPGA 開発キットの BSP に追加できると面白い。

#### 本記事の取り組み
1. Cyclone V GT FPGA 開発キットを使い始める（Lチカ編）
2. Cyclone V GT FPGA 開発キットを使い始める（イーサネット編）
3. 非公式 OpenCL BSP を動かしてみる
4. OpenCL BSP に LED 操作ポートをつけて直接操作する
5. （未完）OpenCL BSP にイーサネットポートをつけて操作する


----

----

# 1. Cyclone V GT FPGA 開発キットを使い始める（Lチカ編）
- 目標 : ボード単独でLチカする

![P_20171203_192133.jpg](https://qiita-image-store.s3.amazonaws.com/0/216085/79c663c8-29a0-b4ae-be88-8db98e482625.jpeg)

- 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 があるのでこれを使えば苦労が無いと思われる

- PATH を通す用の参考スクリプト
 
```shell: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

```cfg:/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"
```

- [サポートページ :arrow_upper_right:](https://www.altera.co.jp/products/boards_and_kits/dev-kits/altera/kit-cyclone-v-gt.html) から Kit Installation をダウンロードして展開
- `cycloneVGT_5cgtfd9ef35_fpga_v13.0.0.1/examples/golden_top/c5gt_pro_goldentop.v` の `define USER` のコメントアウトを解除（モジュール引数最後のコンマは処理）し、以下を追記

```verilog: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`)

```verilog: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`  のコメントアウトを解除（モジュール引数最後のコンマは処理）し、以下を追記

```verilog: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` としてリンクを張る
- [The board support package of Cyclone V GT Development Kit for Intel FPGA SDK OpenCL :arrow_upper_right:](http://www.alterawiki.com/wiki/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` を編集

```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 サンプル :arrow_upper_right:](https://www.altera.co.jp/support/support-resources/design-examples/design-software/opencl/hello-world.html) をダウンロード、展開
- カーネル作成
    + `$ 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） :arrow_upper_right:](https://www.altima.jp/members/p2-faq/altsoftopencl.cfm)
- 結果

```
$ 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によるカメラ制御&画像処理" :arrow_upper_right:](http://www.kumikomi.net/fpga/contents/0014.php) にて DE0-nano-SoC 搭載の Cyclone V SoC 上における例が示されており、全体の参考になる

#### BSP への channel 追加
- c5gt/hardware/c5gt/board_spec.xml に追加 channel を定義

```xml: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 の信号線状態を制御するモジュール

```verilog: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] で以下のようなポート割り当てを実施して保存

![user_led_ctrl.png](https://qiita-image-store.s3.amazonaws.com/0/216085/2e2aef45-735b-9b19-189d-f5c6feba939d.png)

- board.qsys への結線
    + platform designer から board.qsys を開き、user_led_ctrl を追加
    + kernel_clk, kernel_reset を board.qsys 内で配線
    + led_out, kernel_in を export
    + 以下配線例

![user_led_ctrl_wiring.png](https://qiita-image-store.s3.amazonaws.com/0/216085/04522667-7af7-912c-9049-24665208a8cf.png)

- 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 テスト用カーネルコード

```c: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 テスト用ホストコード例

```c: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"));
	}
}
```

- 両者から読まれるインクルードファイル

```c: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.cpp``
    + `CL_CONTEXT_COMPILER_MODE_INTELFPGA=3 ./led_test`

- 結果動画
<blockquote class="twitter-tweet" data-lang="ja"><p lang="ja" dir="ltr">Cyclone V GT Dev Kit の OpenCL BSP から LED を操作できるようにした <a href="https://t.co/0JpmiMH0GY">pic.twitter.com/0JpmiMH0GY</a></p>&mdash; homelith (@homelith) <a href="https://twitter.com/homelith/status/942391442994905088?ref_src=twsrc%5Etfw">2017年12月17日</a></blockquote>
<script async src="https://platform.twitter.com/widgets.js" charset="utf-8"></script>


- ホストから 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 :arrow_upper_right:](https://opencores.org/project,ethernet_tri_mode) （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 への道は遠い