TLDR
OpenCLを使って、HDLで記述した回路のエミュレーションをFPGAでやる方法
Intel FPGA SDK for OpenCLのドキュメント
https://www.altera.co.jp/documentation/mwh1391807965224.html#ewa1452280808662
には、“OpenCL Library"という、VHDL/Verilog等で記述したモジュールをライブラリに変換して、 OpenCLカーネルから呼び出す(実体化)する機能がある。 これを使うことで、自作のIPを容易にFPGAで動作させることができる。
そもそも、FPGAを実用のアプリケーションで使うことが難しいという場合、 色々なレベルでの難しさがある。 以下は、FPGAをGPUのようなアクセラレータとして利用する場合に限定したはなし。
第一に、HDLはプログラミング言語ではないので、それを理解せずに回路を記述するのは「難しい」。
第二に、やりたい応用が数値計算・シミュレーションだとすると、浮動小数点演算を実装する必要があり、 浮動小数点演算回路を作ることが、あるいは、そのためのIPを使うことが「難しい」。
第三に、ある数値計算のアルゴリズムをHDLで設計、記述できたとして、 それを実際にアクセラレータとして使うためには、データを保持するためのメモリも必要だし、 ホストプログラムとのインターフェイスも実装する必要があり、これが「難しい」。
第四に、無事メモリやPCIeのIPコアを自作のHDLで実装した回路と接続できたとして、 ホストプログラムからそれを使うためには、LinuxのデバイスドライバやDMAやらを理解して、 FPGAをPCIe経由で利用するためのコードを実装する必要があり、これが「難しい」。
以上のような多重の困難を解決する1つの方法が高位合成(High Level Synthesis; HLS)と それをサポートするフレームワークになる。 HLSといっても、これまで色々なものが提案されており、それなりに成功したり不成功したりしている。 OpenCLベースのHLSの最も大きな利点は、ホストプログラムとのインターフェイスにOpenCL APIを利用できる、ということにつきる。
OpenCLは、様々なハードウエアで動作することをターゲットとしているので、ホスト側APIは汎用的に定義されている。 そのために、CUDAと比べると色々とまどろっこしいところはあるが、必要なことはやればできる。 GPU用に書いたプログラムは、そのままFPGA用にも転用できる。性能という意味では最適ではないかもしれないけれど。 結果として、OpenCLベースのHLSを利用すると、上記の困難点は全て解決され、 HDLを触る必要もなく、Linuxのカーネル内部やデバイスドライバの仕組みを知ることなく、 あなたの書いたホストプログラムからFPGAを利用できる。最適ではないかもしれないけれど。
HLSでは、ユーザーは普通のプログラミング言語(OpenCLの場合はC言語)で記述すればよい。 HLSのコンパイラがそれ変換して、FPGAで実装できるようなHDLを自動的に生成してくれる。 問題は、この変換はプログラミング言語のコンパイラが機械語を生成する場合と同じく、 一意でもないし、最適な変換を生成することは単純には不可能なため、思ってたのと違う結果になる場合がある。 つまり、HLSのコンパイラのご機嫌を伺って、最適なHDLが生成されるようにプログラムを記述する必要がある。 いわゆる、最適化、というやつ。HLSベースのツールの問題は、このソースコードからHDLへの変換そしてFPGAへ実装、 という一連のサイクルに、1秒でも1分でもなく、1時間のオーダーかかること。 回路規模により1時間から24時間以上まで、様々な場合がある。 コンピューター創世記の、パンチカードを提出して結果が得られるのは半日後、 という時代の利用者の苦労を実感できる、大変教育的な開発環境である。
エミュレーション環境としてのOpenCL for FPGA
本題。HLSを介した最適化のサイクルに1時間とかかかるなら、最初から自分で最適なHDLで実装するほうが効率的かもしれない。 それをOpenCLのフレームワークの中でおこなうための仕組みが"OpenCL Library"になる。
以下、非常に単純化した実例をしめす。VHDLで記述した以下のようなレジスタファイルをOpenCLを介してアクセスしたい。
シングルポートのレジスタファイル mem0.vhd
library ieee;
use ieee.std_logic_1164.all;
use ieee.std_logic_unsigned.all;
entity mem0 is
port (
adr : in std_logic_vector(6 downto 0);
port_a : out std_logic_vector(63 downto 0);
port_c : in std_logic_vector(63 downto 0);
we : in std_logic;
clk : in std_logic;
rst : in std_logic
);
end mem0;
architecture source of mem0 is
subtype REG is std_logic_vector(63 downto 0);
type REG_ARRAY is array (0 to 127) of REG;
signal regfile : REG_ARRAY;
begin
process(clk) begin
if(clk'event and clk='1') then
if (we = '1') then
regfile(conv_integer(adr(6 downto 0))) <= port_c;
port_a <= port_c;
else
port_a <= regfile(conv_integer(adr(6 downto 0)));
end if;
end if;
end process;
end source;
これを実体化するトップレベルのファイルを準備する。
トップレベル chip.vhd
library ieee;
use ieee.std_logic_1164.all;
use ieee.std_logic_unsigned.all;
entity chip is
port (
adr : in std_logic_vector(7 downto 0);
wr : in std_logic_vector(63 downto 0);
we : in std_logic;
rd : out std_logic_vector(63 downto 0);
clk : in std_logic;
rst : in std_logic
);
end chip;
architecture source of chip is
component mem0
port (
adr : in std_logic_vector(6 downto 0);
port_a : out std_logic_vector(63 downto 0);
port_c : in std_logic_vector(63 downto 0);
we : in std_logic;
clk : in std_logic;
rst : in std_logic
);
end component;
signal adr_1 : std_logic_vector(6 downto 0);
signal we_1 : std_logic;
signal wr_1 : std_logic_vector(63 downto 0);
signal rd_2 : std_logic_vector(63 downto 0);
begin
process(clk) begin
if(clk'event and clk='1') then
adr_1 <= adr(6 downto 0);
we_1 <= we;
wr_1 <= wr;
end if;
end process;
mh: mem0 port map (adr => adr_1, port_a => rd_2, port_c => wr_1, we => we_1, clk => clk, rst => rst);
process(clk) begin
if(clk'event and clk='1') then
rd <= rd_2;
end if;
end process;
end source;
“chip.vhd"をOenCL Libraryとして利用するためには、インターフェイスをXMLで記述し、 さらに、今回は自作のIPはVHDLなので、Verilogのラッパーも必要になる。
最初にOpenCLカーネルを示す。
OpenCLカーネル
double pipeline(double a, double b);
__kernel
void
test_pipeline(__global double *data_in_1,
__global double *data_in_2,
__global double *data_out,
const int n
) {
for(int i = 0; i < n; i++) {
double res_from_pipeline = pipeline(data_in_1[i], data_in_2[i]);
data_out[i] = res_from_pipeline;
}
}
このカーネル内の関数「pipeline()」がVHDLの"chip.vhd"と接続される。 OpenCL Libraryでは、カーネル上の変数の種類には意味はなくて、 このカーネル記述は、 “pipeline()は2つの64ビットバス信号を入力とし、1つの64ビットバス信号を出力とするIP"という 事実を意味している。“double"を"long int"に変えても問題ない。 あるいは"double pipeline(double2)“というプロトタイプでも、実質的には同等になる。
今回用意したトップレベルファイルの入力ポートは、7ビットのアドレスと、write enable、64ビットの信号であり、 出力は64ビットの信号になる。OpenCLとVHDLを橋渡しするインターフェースファイルは以下のようになる。
インターフェース定義
<RTL_SPEC>
<FUNCTION name="pipeline" module="v_chip_top">
<ATTRIBUTES>
<IS_STALL_FREE value="yes"/>
<IS_FIXED_LATENCY value="yes"/>
<EXPECTED_LATENCY value="3"/>
<CAPACITY value="1"/>
<HAS_SIDE_EFFECTS value="no"/>
<ALLOW_MERGING value="yes"/>
</ATTRIBUTES>
<INTERFACE>
<AVALON port="clock" type="clock"/>
<AVALON port="resetn" type="resetn"/>
<AVALON port="ivalid" type="ivalid"/>
<AVALON port="iready" type="iready"/>
<AVALON port="ovalid" type="ovalid"/>
<AVALON port="oready" type="oready"/>
<INPUT port="data_in_1" width="64"/>
<INPUT port="data_in_2" width="64"/>
<OUTPUT port="data_out" width="64"/>
</INTERFACE>
<C_MODEL>
<FILE name="RTL/c_model.cl" />
</C_MODEL>
<REQUIREMENTS>
<FILE name="RTL/v_chip_top.v" />
<FILE name="RTL/chip.vhd" />
<FILE name="RTL/mem0.vhd" />
</REQUIREMENTS>
</FUNCTION>
</RTL_SPEC>
「FUNCTION name=“pipeline” module=“v_chip_top”」の部分で、OpenCL側の関数名と、 Verilogのラッパーファイルのモジュール名を対応づける。
“ATTRIBUTES"で重要なのは「EXPECTED_LATENCY value=“3”」の部分。 これは、“chip.vhd"からの出力をラッチするためレイテンシ定義で、これが間違っていると 結果を保持するOpenCL側の配列"data_out[]“の中身がずれる。 今回の回路のばあいは、アドレスを入力してから3サイクル後にデータが出力される。
他、入出力ポートの定義と、必要なHDLファイルの定義の部分を必要に応じて、 変更や追記する。詳しくはドキュメント参照。
Verilogによるラッパー v_chip_top.v
// synopsys translate_off
`timescale 1 ps / 1 ps
// synopsys translate_on
module v_chip_top (
input clock,
input resetn,
input ivalid,
input iready,
output ovalid,
output oready,
input [63:0] data_in_1, // adr 8 , we 1
input [63:0] data_in_2, // wr 64
output [63:0] data_out // rd 64
);
wire [7:0] adr_in;
wire [63:0] wr_in;
wire we_in;
wire [1:0] sel_in;
wire rst_in;
assign ovalid = 1'b1;
assign oready = 1'b1;
assign adr_in = data_in_1[7:0];
assign wr_in = data_in_2;
assign we_in = data_in_1[16];
assign rst_in = resetn;
chip inst(
.adr(adr_in),
.wr(wr_in),
.we(we_in),
.rd(data_out),
.clk(clock),
.rst(rst_int)
);
endmodule
“chip.vhd"を実際にインスタンス化するのはこのファイル。 OpenCL側からくる64bit x 2の信号を、VHDL側の入力ポートに接続している。 アドレスは"data_in_1"の下位8ビット、write enableは"data_in_1"の16ビットとした。 “data_in_2"は、書き込みポートに直結し、“data_out"も読み出しポートに直結する。
あとは"data_in_1[]“と"data_in_2[]“に適切なテストベクトルを送ってやれば、 OpenCL カーネルの定義から、順次VHDL側に信号が送られて、 “data_out[]“で結果を得ることができる。
最後にテストベクトルを設定する、ホストプログラムの一部をしめす。 64ビット整数配列"adrwe[]“を"data_in_1[]“に、 64ビット整数配列"wr[]“を"data_in_2[]“に対応させる場合。
“0xdeadbeef « 32 | i"を先頭から書き込んで読み出す
int c = 0;
// 書き込み
for(int i = 0; i < 8; i++) {
adrwe[c].x[0] = (uint64_t)(((0x1 << 16) | i));
wr[c].x[0] = (uint64_t)0xdeadbeef << 32 | i;
c++;
}
// dummyサイクル
adrwe[c++].x[0] = 0;
adrwe[c++].x[0] = 0;
// 読み出し
for(int i = 0; i < 8; i++) {
adrwe[c++].x[0] = (uint64_t)(i);
}
肝心のOpenCL APIの部分を含めた全体は別途どこかにアップロードする予定。
まとめ
長いです。