情報処理学会研究報告IPSJ SIG Technical Report
OpenCLによるFPGA上の演算と通信を融合した並列処理システムの実装及び性能評価
藤田 典久1小林 諒平1,2山口 佳樹2,1朴 泰祐1,2
概要：近年，高性能計算の分野で再構成可能なハードウェアである Field Programmable Gate Array(FPGA) が次世代の演算加速装置として注目されている．FPGA を高性能計算で用いる際の障壁は開発の困難さであったが，高位合成手法の発展に伴いこの問題は解決しつつある．最新の FPGA は最大で100Gbps×4 の通信性能を有しており，我々はその強力な通信性能に注目している．FPGA の絶対性能は
他のアクセラレータよりも低いが，FPGA が持つ演算能力と通信能力を組み合わせることでより広い範囲の問題に FPGA が適用できると考えている．本研究の目的は，高位合成で記述された FPGA アプリケーションから通信機構を操作し並列処理システムを実現することである．通信のスループットやレイテンシだけでなく，姫野ベンチマークを用いた性能評価を行い，高位合成で記述した FPGA アプリケーションで並列計算が可能なことを示す．我々は FPGA 間で直接通信を行う環境として Channel over Ethernet (CoE)
というシステムを開発しており，バンド幅は最大で 7.13Gbps を達成し，4 バイト通信時のレイテンシは980ns であった．姫野ベンチマークで，問題サイズ M を 4 FPGA で実行する場合に 22659 MFLOPS の性能が得られ，4 FPGA 時に 1 FPGA 時と比べて 3.61 倍という良好な Strong Scaling の結果が得られた．
1.はじめに
近年，高性能計算の分野で再構成可能なハードウェアで
ある Field Programmable Gate Array (FPGA) が次世代
の演算加速装置として注目されている．高性能計算では絶
対性能に加えて電力効率を高めることが重要な問題であ
り，FPGA を用いてこの問題を解決しようとしている．
従前，FPGA を高性能計算で用いる際の障壁は開発の
困難さであったが，高位合成 (High Level Synthesis; HLS)
手法の発展に伴いこの問題は解決しつつある．FPGA 開発
では，Verilog HDL や VHDL といった ハードウェア記述
言語 (Hardware Description Language; HDL) を用いて回
路を記述することが一般的である．HDL は 1 クロック・
1 ビット単位で回路の動作を記述するもので，利用に際し
てハードウェアの知識が必要となり，計算科学者が FPGA
を用いる際の障壁となっていた．一方，高位合成は C や
C++等のソフトウェア開発で用いられる言語を用いてハー
ドウェアを記述するものであり，ハードウェア開発に精通
していない科学者が FPGA を利用することを可能とする
ものである．
現在の FPGA の絶対性能は，アクセラレータとして広
く用いられている Graphics Processing Unit (GPU) に敵
1筑波大学 計算科学研究センター2筑波大学 システム情報工学研究科
わない．したがって，既に GPU で高速に計算できる問題
を FPGA に適用しても，GPU の性能を超えられるとは言
い難い．FPGA を高性能計算で用いる場合は，アプリケー
ションの中のどの計算が FPGA に適するかを見極めて適
用することが重要だと言える．
我々は FPGA の持つ強力な通信性能に注目している．
最新の FPGA は最大で 100Gbps×4 の通信性能を有してい
る．また，それらの通信機構は直接 FPGA に接続されてお
り，オーバーヘッドの少ない FPGA 間通信を可能とする．
NVIDIA GPU では GPUDirect for RDMA (GDR) [1] と
呼ばれる技術があり，外部の通信機構が GPU のメモリに
直接アクセスができる．しかしながら，通信の制御などは
CPU にあり，GPU が主体となって通信できる技術ではな
く，CPU-GPU 間の同期コストや PCI Express (PCIe) を
経由することによるオーバーヘッドは残っている．FPGA
を単体のアクセラレータとして見ると適用できる範囲は小
さいが，FPGA が持つ演算能力と通信能力を組み合わせる
ことでより広い範囲の問題に FPGA が適用できると考え
本研究の目的は，高位合成で記述された FPGA アプリ
ケーションから通信機構を操作し並列処理システムを実現
することである．高位合成から通信できるシステムを開発
し，それの性能評価を行う．通信のスループットやレイテ
c⃝ 2018 Information Processing Society of Japan1
Vol.2018-HPC-167 No.9
2018/12/17
情報処理学会研究報告IPSJ SIG Technical Report
ンシだけでなく，姫野ベンチマークを用いた性能評価を行
い，高位合成で記述した FPGA アプリケーションで並列
計算が可能なことを示す．
2.関連研究
OpenCL を FPGA で用いてアプリケーションやベンチ
マークの性能評価を行った論文はいくつか報告されてい
る．[2] では，元々 GPU 向けに作成されたコードをその
まま FPGA 向けに用いても性能が悪く，OpenCL コード
が FPGA 向けに最適化されている必要があると述べられ
ている．HPC 研究会においても，[3] や [4] で FPGA と
OpenCL を用いた研究報告がなされているが，どちらでも
OpenCL の最適化の困難さ，すなわち，CPU や GPU と異
なる記述スタイルが必要であると述べられている．FPGA
の絶対性能は GPU など他のアクセラレータと比べると低
く，どのような種類の処理を FPGA にオフロードするの
かが重要となる．
本研究の独自性は，高位合成が持つ高い生産性と，FPGA
が持つ高い通信能力を併せて利用し，並列計算を行うこと
である．FPGA と OpenCL を用いる研究は行われている
が 1 FPGA のみを用いているものであり，並列計算が行わ
れている例は知られていない．
3.Intel FPGA SDK for OpenCL
Intel 社は自社の FPGA 向けの高位合成の処理系として，
Intel FPGA SDK for OpenCL を公開している．この SDK
を用いることで，OpenCL 言語を用いて FPGA をプログ
ラミングできる．OpenCL から FPGA 上のハードウェア
を生成するだけでなく，ホスト CPU で用いるドライバや
ランタイムライブラリも提供されており，この SDK のみ
でシステム全体を構築できる．
OpenCL 利用時の FPGA ハードウェアの構造の概略図
を図 1 に示す．図からわかるように，FPGA の内部構造
は，大きくわけて 2 つの部分から構成される．1 つは Board
Support Package (BSP) 由来の部分，もう一つはコンパイ
ル対象の OpenCL コードから由来する部分である．
BSP は本 OpenCL SDK 固有の要素であり，異なる FPGA
ボード上で同じ OpenCL プログラムを扱うために存在す
る．OpenCL 対応の FPGA ボードは多数あり，各ボード
でハードウェアの構成が異なる．例えば，それぞれのボー
ドで搭載されているメモリの種類や FPGA チップが異な
る場合がある．したがって，それぞれのボードに固有の情
報を OpenCL コンパイラに与える必要があり，BSP がそ
の役割を担う．BSP にはペリフェラルコントローラが含ま
れており，OpenCL のアクセラレータとして利用するため
の最低限のものとして，PCIe コントローラと Dual Data
Rate (DDR) メモリコントローラが含まれる．
FPGA Chip
BSPOpenCL Kernels 
a.clb.clc.cl
Partial Reconﬁguration 
図 1: OpenCL 利用時における FPGA の内部構造の概略
3.2通信機構に対応する BSP
通信機構を持つ OpenCL 対応の FPGA ボードであって
も，そのボード用の BSP にコントローラが含まれている
ことは一般的ではなく，OpenCL から通信機構を操作し
並列計算を行うためには，制御用のロジックを BSP に組
み込む必要がある．我々は，これまでの研究 [5], [6] で，
OpenCL の BSP に通信機構のコントローラを追加し，そ
れを OpenCL コードから制御できることを明らかにして
本稿で実装した OpenCL 通信機構は，これらの研究で得
られた知見を元に開発したものである．なお，BSP に対し
てコントローラを追加し，OpenCL カーネルから用いる手
法の詳細については本稿では省略する．前述した研究会報
告と論文を参照して頂きたい．
3.3Channel 拡張
Intel FPGA SDK for OpenCL は OpenCL 言語に対して
いくつかの FPGA に特化した拡張を加えている．拡張の 1
つに，“Channel” 機構がある．Channel は OpenCL カーネ
ル間でデータを直接交換するパイプのようなものである．
カーネル間でデータをやりとりする場合，グローバルメ
モリに経由してデータを交換する方法が一般的であるが，
Channel を用いる場合は FPGA 内部にデータパスが構成
され，チップ外にあるメモリにアクセスすることなく通信
ができる．したがって，従来手法と比べて高性能であり，
メモリアクセスに関する回路が生成されないため省リソー
スとなる．Channel には “I/O Channel” と呼ばれる種類の
Channel があり，これはカーネル間ではなく，BSP とカー
ネルの間を接続するために存在し，主に BSP にあるペリ
フェラルコントローラと OpenCL カーネル間の接続に用
本研究では，図 2 にある様に，BSP に通信機構を操作
するためのコントローラ (40Gbit Ethernet) を追加し，そ
れを OpenCL コードから操作するために I/O Channel を
用いる．また，Channel を用いて複数のカーネルの間を接
続し，通信に必要な処理を実現する．通信機構の実装の詳
c⃝ 2018 Information Processing Society of Japan2
Vol.2018-HPC-167 No.9
2018/12/17
情報処理学会研究報告IPSJ SIG Technical Report
OpenCL Kernels 
40G Eth. Controller 
BSPOpenCL Circuits
Serial Links (x4)IO Channels
図 2: BSP に含まれる通信機構のコントローラと OpenCL
カーネル間の接続．
__kernel void sender(__global float* restrict x, int n) {
for (int i = 0; i < n; i++) {
float v = x[i];write_channel_intel(simple_out, v);
__kernel void receiver(__global float* restrict x, int n) {
for (int i = 0; i < n; i++) {
float v = read_channel_intel(simple_in);x[i] = v;
sender code on FPGA1
receiver code on FPGA2
図 3: CoE を用いて通信を行うコード例．
細については第 4 章で述べる．
4.CoE: Channel over Ethernet
我々は FPGA 間で直接通信を行う環境として Chan-
nel over Ethernet (CoE) というシステムを開発している．
CoE の基本コンセプトは，前述した Channel による通信
を FPGA 内だけでなく，ノード間に拡張するものである．
CoE システムは，BSP に追加された Ethernet コントロー
ラおよび周辺回路 (HDL で記述) ，OpenCL カーネルで記
述された制御カーネル群と，それらの間をつなぐ Channel，
CoE システムとアプリケーションの境界にある Channel
である “CoE Channel” から構成される．
図 3 に CoE を用いて通信する際の簡単な例を示す．2 つ
のカーネル関数は異なる FPGA 上で動作しており，送信側
の “simple out” channel と受信側の “simple in” channel
が CoE システムを通じて繋っており，通信が可能となる．
CoE の通信プロトコルはその名前が示す通り Ethernet
を用いている．FPGA で利用できるプロトコルは Ethernet
だけではないが，その中から Ethernet を採用している理
由は，市販されているスイッチを用いて多数の FPGA か
らなるネットワークを構築できるからである．
Ethernet を用いるということは，Ethernet の仕様に従っ
たパケットフォーマットに従って通信をする必要がある
が，CoE ではそれらを行う回路を OpenCL 側で実装して
いる．この方式のメリットは OpenCL 側と BSP 側のイン
ターフェイスが送信・受信のあわせて 2 つの I/O Channel
で固定される点にある．もし，パケット構築などを BSP
側に構築してしまうと，アプリケーション側で使われる
Channel インターフェイスが変化するたびに BSP を再構
築しなければならない．
HDL は生産性が低く機能開発や動作確認に時間がかかる
こと，外部インターフェイスで要求される動作周波数*1を
満すためのチューニングの手間が大きい，といった点から
BSP の変更はコストが高いため留めたい．一方で，こう
いった処理を OpenCL で実装することは性能面での低下が
懸念される．通信スループット性能については外部リンク
に性能の上限があるため，OpenCL で実装したとしても十
分満せると考えられるが， OpenCL ではサイクルレベルで
の動作記述ができないため，レイテンシの最適化が困難で
あること，一般的に高位合成は HDL で記述する場合より
も多くの回路リソースを消費するという問題点がある．こ
れらの要素を考慮して，制御に関するロジックを OpenCL
で記述するメリットの方が大きいと判断し，OpenCL で記
述する方針を選択している．
4.2一対一通信
CoE における一対一通信の送信側のデータの流れを図 4
に示す．この図は姫野ベンチマークにおける CoE ネット
ワークを表したものであり，4 つの袖領域交換用 (X 次元
と Y 次元用にそれぞれ 2 つ) の channel と，Allreduce の
ための channel がある．
CoE の送信回路は以下に示す 4 つのカーネルから構築さ
れており，全て OpenCL で記述されている．
expanderカーネルからの入力データを 128bit 単位に
パッキングして buﬀer に送り出す．
buﬀerexpander からのデータをバッファリングし，Eth-
ernet パケットに整形しパケットストリームを作成
scaler単一の値を含むパケットを生成する．バッファリ
ングをしないため高速である．
multiplexer複数のパケット列をマージし，1 つのスト
リームを形成する．
CoE は 2 つの送信モードを持つ．1 つはバッファリング
モードであり，高いスループットを得るためにある程度の
データをバッファリングし，1 つのパケットに纏めて送信
する．このモードでは expander カーネルと buﬀer カーネ
ルを用いてパケットが生成される．もう 1 つはスカラー
モードであり，これは単一の値を低レイテンシで送信する
ことを意図している．このモードは集団通信を実装する際
に用いることを想定しており，バッファリングを行わない
ため低リソース・低レイテンシな通信を行える．このモー
ドでは scalar カーネルを用いてパケットが生成される．
受信側は複雑なバッファ処理が必要ないため，送信側よ
りシンプルに構成されている．パケットに含まれている宛
先 Channel の ID から経路を判断して Channel にデータ
*1例えば，PCIe バスに関する回路は 250MHz，40GbE に関する回路は 312.5MHz で動作することが求められる．
c⃝ 2018 Information Processing Society of Japan3
Vol.2018-HPC-167 No.9
2018/12/17
情報処理学会研究報告IPSJ SIG Technical Report
を分別するカーネルと，し Ethernet ヘッダの除去を行う
カーネルの 2 つから構成されており，単一の動作モードか
CoE Channel はそれぞれ固有の ID を持っており，ID と
宛先アドレスで通信相手を決定する．宛先は送信用パケッ
トを生成するカーネルを起動する際に引数で指定する．ホ
スト側のコードと OpenCL カーネルの間のインターフェイ
スをカーネル起動で行うため，カーネルの実行モデルに由
来する制限がある．OpenCL の仕様では実行中のカーネル
をホストから停止する手段がない．すなわち，一度カーネ
ルを起動して宛先を決定すると，プログラム実行中は変更
できない．しかしながら，高性能計算のアプリケーション
における一対一通信は通信相手が固定であることが一般的
であるため，この制限はさほど問題にならないと考えてい
る．また，Intel OpenCL 環境には Host Channel と呼ばれ
るホストと FPGA 間を Channel 接続する機能があり，こ
の機能を使えば実行中のカーネルにデータを渡せるため，
この制限を解消できるものと考えており，今後実装する予
4.3集団通信: Allreduce
CoE は一対一通信だけでなく集団通信も同様のアプロー
チで実装している．現時点で CoE に実装されている集団
通信は Allreduce のみであるため，本節では Allreduce に
ついてのみ述べる．
図 5 に CoE で用いている Allreduce のアルゴリズムを
示す．ツリー型のアルゴリズムを採用しており，それぞれ
の正方形はカーネルを表している．“+” は入力を加算して
出力するカーネル (add)，“=” は入力をそのまま出力にコ
ピーするカーネル (dup) である．このアルゴリズムは一対
一通信のみで構築できるため CoE のシステムに適するこ
とと，演算順序が固定であり浮動小数点数の演算を行なっ
ても最終的に全てのプロセスで同じ結果が得られる点が利
前節で述べた通り，CoE の一対一通信は通信相手を高
頻度に切り替えることを想定していない．したがって，一
対一通信で構築される Allreduce もそれに由来する制限が
あり，Allreduce に参加する最大ノード数を予め決定しツ
リーの段数を静的に決めなければならない．add も dup も
2 入力 2 出力 channel のカーネルであり，カーネル引数で
入力 channel と出力 channel の有効無効を切り替えられる．
ノード数が最大数よりも少ない場合は，ツリー中に何もし
ない段を作ることで同じ回路で Allreduce が可能である．
現在の実装で対応しているデータ型は ﬂoat のみ，対応し
ている演算も MPI SUM 相当のみであるが，演算ロジック
が OpenCL で組まれているため，他の演算やデータ型の追
加は容易であると考えている．
4.4制限事項
CoE は開発中のシステムであり，いくつかの制限事項が
残っている．
• アプリケーションと CoE システムの間の Channel は
ﬂoat 型しか取れない．
• multiplexer におけるバッファリングの不足により，同
時に 1 つの channel しか通信できない．
• フロー制御や再送制御の実装はなく，受信側のバッ
ファが不足した場合はパケットが脱落する．
特にバッファリングの機能不足により通信性能が制限
されている．複数の channel に対して ﬂoat 型のデータを
書き込んでも，あるクロックサイクルで送信されるデー
タは 1 つの channel からしか選ばれない．言い換えると，
32bit/cycle の通信能力しか持たないということである．こ
れらの制限事項は引き続き開発を行い解消してく予定で
5.姫野ベンチマークの CoE 実装
5.1FPGA 実装および最適化
姫野ベンチマークは非圧縮流体解析を行うベンチマーク
プログラムである [7]．姫野ベンチマークの計算のパター
ンは高性能計算で典型的なステンシル計算を行うもので
あり，ポアソン方程式をヤコビ反復法を用いて解く．姫野
ベンチマークは C および Fortran で記述されており CPU
向けの実装はあるが，FPGA 向けの実装は存在しないた
め，まずは FPGA 向けの実装を行う必要がある．本研究で
は，C 言語版の姫野ベンチマークをベースとして，FPGA
OpenCL 版の実装を行う．
姫野ベンチマークは配列 p に関して 19 点ステンシル計
算であり，配列 p を複数回参照するコードである．通常，
CPU や GPU などのプロセッサで実行する場合はキャッ
シュを通じてメモリにアクセスをするため，配列 p に対す
るアクセスはキャッシュヒットすることが期待でき，配列
p のアクセスで消費されるメモリ帯域が減少する．しかし
ながら，現在の OpenCL 環境ではメモリアクセスに対す
るキャッシュの作成は限定的であり，ステンシル計算のパ
ターンではキャッシュがないに等しい．
FPGA でステンシル計算のメモリアクセスを最適化する
場合，シフトレジスタを用いて最適化する手法が一般的で
ある [8]．また，同手法が OpenCL で記述する際にも有効
なことが明らかとなっている [2]．本研究における姫野ベ
ンチマークの実装も，同様の手法で最適化を行っている．
DDR メモリから読み出した配列 p の要素を再利用のため
にシフトレジスタに格納し，再び同じ要素にアクセスす
る際は DDR メモリからではなくシフトレジスタから取り
c⃝ 2018 Information Processing Society of Japan4
Vol.2018-HPC-167 No.9
2018/12/17
情報処理学会研究報告IPSJ SIG Technical Report
expander buffer 
expander buffer 
expander buffer 
expander buffer 
scalar scalar scalar scalar 
binary tree merge network
Allreduce
Himeno Kernels 
[stage1][stage2][stage3][stage4]
packet (280bit) packed data (136bit) input data (32bit) 
図 4: CoE 送信回路の構成．
p0p1p2p3
p0p1p2p3
図 5: Allreduce 実装の通信パターン図．p0～p4 はそれぞ
れプロセス，赤矢印は FPGA 間通信，黒矢印は FPGA 内
の通信を表す．
5.2CoE の適用
姫野ベンチマークの CPU 向け C 言語版のカーネル部全
体の疑似コードを図 6 に示す．カーネルは大きくわけて，
演算部 (A)，更新部 (B)，通信部 (C) の 3 部分から構成さ
れている．CPU の実装では，ノード間の並列化の通信手法
に Message Passing Interface (MPI) を用いており，通信
部で MPI 関数を呼び出すことで袖領域の更新および gosa
変数のリダクションを行う．
ステンシル計算では，通信最適化の 1 つとして，通信
と計算のオーバーラップがよく用いられる [9]．オーバー
ラップを行う際は，各タイムステップでまず袖領域の値を
先に計算し，その後，袖領域の通信とそれ以外の領域の計
算を同時に行うことで，通信時間の隠蔽を図る．姫野ベン
チマークでは，計算と通信はオーバーラップしておらず，
計算 (A+B) が終ってから袖領域の通信 (C) を行うが，通
信と計算のオーバーラップを適用可能であり，適用しても
計算は破綻しない．
図 7 に CoE を適用した姫野ベンチマークコードの一部
を示す．この部分は図 6 の A の部分のループに相当する
部分である．write channel intel 関数が 4 つ記述され
ているが，この Channel への書き込みによってデータが
for (i = 1; i < imax-1; i++)
for (j = 1; j < jmax-1; j++)
for (k = 1; k < kmax-1; k++) {
s0 = ...;ss = ...;wgosa += ss * ss;wrk2[i][j][k] = p[i][j][k] + omega * ss;
for (i = 1; i < imax-1; i++)
for (j = 1; j < jmax-1; j++)
for (k = 1; k < kmax-1; k++) {
p[i][j][k] = wrk2 [i][j][k];
sendp(ndx, ndy, ndz);MPI_Allreduce(&wgosa, &gosa, ...);
図 6: CPU 版の姫野ベンチマークの計算の流れ．
CoE システムに渡り，最終的に隣接 FPGA の Channel に
配送される．“if (enable)” で囲まれた演算部と，4 つの
write channel intel がシーケンシャルに実行されるよう
に見えるが，FPGA 上ではパイプラインのハードウェアと
して表現されているため，これらの要素はそれぞれ並列し
て動作する．
姫野ベンチマークに CoE を適用する場合を考えると，
CoE の通信はパイプラインの動作モデルを前提しているた
め，MPI のように通信と計算のオーバーラップを行うとい
う概念が不要なことがわかる．計算と通信が一体のパイプ
ラインを形成しているため，計算を司る回路と通信を司る
回路は平行に動作し，計算をしながら通信を行うという挙
動が自然に記述できる．
6.性能評価
6.1評価環境
性能評価には Pre-PACS version X (PPX) クラスタシス
テムを用いる．PPX は筑波大学 計算科学研究センターで
運用中のシステムであり，同センターが開発を計画してい
る PACS シリーズ・スーパーコンピュータ次世代機のプロ
トタイプシステムである．
PPX は Intel FPGA を持つノード，Xilinx FPGA を持
c⃝ 2018 Information Processing Society of Japan5
Vol.2018-HPC-167 No.9
2018/12/17
情報処理学会研究報告IPSJ SIG Technical Report
bool enable = (1 <= i && i < t.imax - 1) &&
(1 <= j && j < t.jmax - 1) &&(1 <= k && k < t.kmax - 1);
float value = 0;if (enable) {
s0 = ...;ss = ( s0 * v.a[3] - v.p[18] ) * v.bnd;wgosa += ss*ss;value = v.p[18] + t.omega * ss;t.wrk2[iter] = value;
}if (!t.last && t.pi > 0 && enable && i == 1) {
write_channel_intel(himeno_comp3_i_neg, value);
}if (!t.last && t.pi < t.ndx0 - 1 &&
enable && i == t.imax - 2) {
write_channel_intel(himeno_comp3_i_pos, value);
}if (!t.last && t.pj > 0 &&
enable && j == 1) {
write_channel_intel(himeno_comp3_j_neg, value);
}if (!t.last && t.pj < t.ndy0 - 1 &&
enable && j == t.jmax - 2) {
write_channel_intel(himeno_comp3_j_pos, value);
図 7: CoE を適用した姫野ベンチマークにおける袖通信部
つノードの 2 グループからなるが，本研究では Intel FPGA
のみを持ちいる．表 1 に PPX システムの Intel FPGA
ノードの仕様を示す．各ノードに Broadwell Xeon CPU
×2, NVIDIA P100 GPU×2, InﬁniBand EDR HCA×1,
BittWare FPGA ボード ×1 が搭載されている．また，図 8
にあるように CPU, GPU 向けの InﬁniBand ネットワーク
だけでなく，FPGA だけが接続されている専用ネットワー
クがあり Mellanox 社製のスイッチを用いて FPGA 間が接
続されている．
本稿の性能評価は最大で PPX 4 ノードを用いて行う．
スイッチはポートあたり最大で 100Gbps の能力を有する
が，FPGA ボード側が最大で 40Gbps までの通信しか対応
していないため，40Gbps の速度で用いる．また，それぞ
れの FPGA ボードは 2 つの QSFP+ポートを持つが，今回
の実験では片側のみ利用している．
6.2pingpong ベンチマーク
CoE 通信機構の基礎性能を評価するために pingpong ベ
ンチマークの評価を行う．pingpong ベンチマークは図 9
の様に実装しており，2 つの FPGA 間でデータを交換し
性能を測定する．ただし，2 つの FPGA 間は直結ではなく
スイッチを用いて接続しているため，それによるオーベー
ヘッドが含まれている．
時間の測定は図 10 にあるように，3 箇所を OpenCL カー
ネルの動作クロック精度で時間を測定している．Arria 10
FPGA で OpenCL を利用する場合，カーネルコードによっ
表 1: 評価環境
CPUIntel Xeon E5-2660 v4 × 2
CPU MemoryDDR4 2400 MHz 64 GB
(8 GB × 8)
Host OSCentOS 7.3
Host Compilergcc 4.8.5
OpenCL SDKIntel FPGA SDK
for OpenCL 17.1.2.304
FPGABittWare A10PL4
(10AX115N3F40E2SG)
FPGA MemoryDDR4 2133 MHz 8 GB
(4 GB × 2)
Communication PortQSFP+ × 2
(40 Gbps × 2)
Ethernet SwitchMellanox MSN2100
図 8: PPX のネットワーク図．
kernel pingkernel reply
kernel pong 
FPGA1FPGA2
図 9: pingpong ベンチマークの構成．
ping send
pong recv
t0 t1  t2
図 10: pingpong ベンチマークにおける時間計測．
て動作周波数は変化するが，一般的に 200～250MHz の範
囲になり，4～5ns の精度で時間を測定できる．通信レイテ
ンシおよびバンド幅は片道レイテンシに通信時間を足したt1 − t0
2+ (t2 − t1)[s] を元に計算する．
レイテンシの測定結果を図 11 に，バンド幅の測定結果を
図 12 に示す．なお，今回測定に使用した実装は 230MHz
c⃝ 2018 Information Processing Society of Japan6
Vol.2018-HPC-167 No.9
2018/12/17
情報処理学会研究報告IPSJ SIG Technical Report
48163264128256
Data Size [B]
Latency [ns]
Latency Benchmark Result (4B - 2MB)
図 11: レイテンシの測定結果．
101102103104105106
Data Size [B]
Bandwidth [Gbps]
Throughput Benchmark Result (4B - 2MB)
図 12: バンド幅の測定結果．
で動作しており，1 クロックサイクルの時間は 4.35ns であ
る．Ethernet のパケット最大サイズは 512byte になるよう
に作成する．ただし，Ethernet ヘッダ (12 バイト) および
CoE ヘッダ (2 バイト) でパケットあたり 16 バイトを消費
するため，ペイロードは 1 パケットあたり最大で 496 バイ
CoE 通信のバンド幅は最大で 7.13Gbps を達成し，CoE
を用いた 4 バイト通信時の最小レイテンシは 980ns であっ
た．ただし，OpenCL カーネルのクロックに同期して動く
ため，レイテンシは OpenCL カーネルの動作周波数に影響
を受ける．また，送信回路のサイクル数がアプリケーショ
ンからいくつの Channel が CoE に接続されているか依存
し，アプリケーションの構造によってもレイテンシは変動
するため，全てのアプリケーションでこの値になるという
ことではない．しかしながら，FPGA による直接通信が低
レイテンシで行えていることを十分示せているといえる．
6.3姫野ベンチマーク
本節では CoE を適用した姫野ベンチマークの性能を示
す．第 5 章で述べた通り，姫野ベンチマークの実装はシフ
トレジスタを用いて最適化している．また，袖領域通信に
必要なバッファを FPGA 内部のメモリに格納しており，こ
Number of FPGAs
Performance [M FLOPS]
図 13: 姫野ベンチマークの性能のグラフ．
表 2: 姫野ベンチマークの性能 [M FLOPS]．
1 FPGA2 FPGAs4 FPGAs
XS5,0219,86116,633
S6,10711,33621,109
M6,27312,34522,659
表 3: 各問題サイズで 1 FPGA を 1 とした場合の速度向上
1 FPGA2 FPGAs4 FPGAs
XS1.001.963.31
S1.001.863.46
M1.001.973.61
れらのメモリ使用量は 1 FPGA あたりの問題サイズに比
例する．そのため，本稿で用いた FPGA では M サイズま
での問題しか解くことができない．
オリジナルの姫野ベンチマークは実行時間が 1 秒になる
ように反復回数を決定し性能を測定するが，それでは問題
サイズ，FPGA 数，実行毎の性能ゆらぎなどの影響で反復
回数が変化してしまい性能比較が難しい．そのため本稿で
は反復回数を固定して測定しており，それぞれ XS=10000
回，S=1250 回，M=150 回に固定している．また，利用す
る FPGA 数を 1, 2, 4 FPGA と変化させ，ストロングス
ケーリングの条件で測定を行う．
CoE を適用した姫野ベンチマークの性能を図 13 と表 2
に示す．また，カーネルの動作周波数を表 4 に，それぞれ
の問題サイズで 1 FPGA の時の性能を 1 とした性能向上
率を表 3 に示す．問題サイズ M を 4 FPGA で実行する場
合に 22659 MFLOPS の性能が得られ，また，問題サイズ
M で 4 FPGA 時に 1 FPGA 時と比べて 3.61 倍という良好
なストロングスケーリングの結果が得られた．
CoE 通信のバンド幅は最大で 7.13Gbps を達成したが，
物理層として 40GbE を用いていることを鑑みるとこの性能
c⃝ 2018 Information Processing Society of Japan7
Vol.2018-HPC-167 No.9
2018/12/17
情報処理学会研究報告IPSJ SIG Technical Report
表 4: 姫野ベンチマークのカーネル動作周波数 [MHz]．
1 FPGA2 FPGAs4 FPGAs
XS210.00227.08227.08
S226.19219.44220.37
M219.44220.83210.42
表 5: 姫野ベンチマーク XS サイズにおける問題サイズと
メモリサイズの比較．
XS 1 FPGAXS 2 FPGAsXS 4 FPGAs
メモリサイズ(33,33,65)(19,33,65)(19,19,65)
内点サイズ(30,30,62)(15,30,62)(15,15,62)
FLOP/mesh26.8023.2820.21
FLOP/mesh5628.510474.116817.3
ratio×1.00×1.88×3.26
は十分ではない．この性能の差は，1 クロックあたり 32bit
データしか通信できないという現在の実装の制限から来る
ものである．ベンチマークの回路が 230MHz で動作してい
るため，得られる最大の性能は 0.23 × 32 = 7.36Gbps であ
る．今回の実装のピークに対して 96%の性能が得られてお
り，システムは想定通り動作しているといえるが，スルー
プットの改善が今後の課題である．
姫野ベンチマークの性能測定結果から，問題サイズ XS
の 4 FPGA の時のスケーリングが 3.3 倍となっていること
がわかる．この場合並列化効率は 83%であり，FPGA の低
レイテンシな通信を活かせていないように見える．しかし
ながら，この問題の原因は，通信ではなく，Strong Scaling
により問題サイズが小さくなることで，内点と袖領域の
比が悪化していることにある．現在の姫野ベンチマーク
の実装はメッシュデータのサイズ (パラメータ MIMAX,
MJMAX, MKMAX) におおむね比例し*2，動作周波数に
よっても性能が異なる．メッシュサイズと内点サイズから
求めた 1 メッシュあたりの FLOP と，それに対して動作
周波数を掛けた値の比較を表 5 に示す．表 5 のデータよ
り，XS サイズで並列化効率が悪い原因は，通信によるオー
バーヘッドではなく，問題サイズに対して袖領域の占める
割合が多いことであるとわかる．
図 6 からわかるように，姫野ベンチマークは A で計算を行
い wrk2 配列に書き込み，B でメモリコピーを行い配列 p に
新しいメッシュデータを書き込む．FPGA 実装において B
の部分は内部メモリバス幅と同じ 512bit 幅でコピーを行う．
したがって，XS 4 FPGA の最も 1 FPGA あたりの問題サイ
ズが小さい場合でも，少なくとも 15 × 15 × 62
16= 871.875
クロックの時間が必要であり，実際には，これにメモリレ
イテンシやオーバーヘッドなどが加わるため，必要な時間
はさらに伸びると考えられる．CoE の通信レイテンシは
*2姫野ベンチマークは計算に必要な袖領域に加えてパディングを加えるため，内点サイズ+2 よりもメモリサイズは大きくなる．
pingpong ベンチマークの結果より 1µs 程度であり，この
時間は 220 クロック (220MHz 動作時) に相当する．した
がって，B のメモリコピーにかかる時間の方が通信にかか
る時間よりも長く，通信がボトルネックになっていないと
考えられる．
5.2 節で，CoE においては計算と通信が一体のパイプラ
インを形成しているため，MPI のように通信と計算をオー
バーラップさせるという概念が不要であり，計算をしなが
ら通信を行うという挙動が自然に記述できると述べた．前
述したメモリコピーと通信時間の比較は，計算と通信がパ
イプラインでほぼ同時に進行するからできることであり，
CoE の通信モデルは FPGA で並列計算を行う際に適して
いると考えられる．
8.おわりに
本研究では，Intel OpenCL 開発環境が FPGA 向け拡張
として実装している Channel API を，異なる FPGA 間に
拡張するコンセプトである CoE 環境を開発し，性能評価
を行なった．システムの性能を評価するために pingpong
ベンチマークと姫野ベンチマークに CoE を適用した．
CoE 通信のバンド幅は最大で 7.13Gbps を達成し，CoE
を用いた 4 バイト通信時の最小レイテンシは 980ns であっ
た．姫野ベンチマークで，問題サイズ M を 4 FPGA で実
行する場合に 22659 MFLOPS の性能が得られ，また，問
題サイズ M で 4 FPGA 時に 1 FPGA 時と比べて 3.61 倍
という良好なストロングスケーリングの結果が得られた．
しかしながら，さらなる最適化の必要や，フロー制御なエ
ラー検出など通信機構として実用するのに必要な機能がま
だ不足しており，今後開発を進めていく予定である．
CoE においては計算と通信が一体のパイプラインを形成
しているため，計算をしながら通信を行うという挙動が自
然に記述でき，通信隠蔽を行いやすい．したがって，CoE
の通信モデルは FPGA で並列計算を行う際に適している
と考えられる．我々は宇宙物理学のアプリケーションの
FPGA 向け最適化も行なっており [10]，今後，CoE の通信
システムをそのアプリケーションに適用し，複数 FPGA を
用いた並列計算を行う予定である．
また，筑波大学 計算科学研究センターでは，次期スー
パーコンピュータ (名称: Cygnus) を 2019 年 5 月から運用
する．Cygnus は 1 ノードに 2 CPU, 4 GPU, 2 FPGA を
持つマルチヘテロジニアスなシステムとなり，FPGA 間
は 4×100Gbps の通信を用いた 2D トーラスネットワーク
となる．今後は本システムを Cygnus 上で稼動させるため
に，100Gbps 通信対応および 2D トーラスネットワーク対
応を行う予定である．
謝辞本研究の一部は，「高性能汎用計算機高度利用事
業」における課題「次世代演算通信融合型スーパーコン
ピュータの開発」及び，文部科学省研究予算「次世代計算
c⃝ 2018 Information Processing Society of Japan8
Vol.2018-HPC-167 No.9
2018/12/17
情報処理学会研究報告IPSJ SIG Technical Report
技術開拓による学際計算科学連携拠点の創出」による．ま
た，本研究の一部は，「Intel University Program」を通じ
てハードウェアおよびソフトウェアの提供を受けており，
Intel の支援に謝意を表する．
[1]NVIDIA Corporation:GPUDirect for RDMA, (on-line),availablefrom⟨https://docs.nvidia.com/cuda/gpudirect-rdma/index.html⟩.
[2]Zohouri, H. R., Maruyama, N., Smith, A., Matsuda,M. and Matsuoka, S.:Evaluating and OptimizingOpenCL Kernels for High Performance Computingwith FPGAs, Proceedings of the International Confer-ence for High Performance Computing, Networking,Storage and Analysis, SC ’16, Piscataway, NJ, USA,IEEE Press, pp. 35:1–35:12 (online), available from⟨http://dl.acm.org/citation.cfm?id=3014904.3014951⟩(2016).
[3]大島聡史，塙敏博，片桐孝洋，中島研吾： FPGA を用いた疎行列数値計算の性能評価，情報処理学会研究報告,2016-HPC-153 (2016).
[4]塙敏博，伊田明弘，大島聡史，河合直聡： FPGA を用いた階層型行列ベクトル積，情報処理学会研究報告,2016-HPC-155 (2016).
[5]大畠佑真，小林諒平，藤田典久，山口佳樹，朴泰祐： OpenCL と Verilog HDL の混合記述による FPGA間 Ethernet 接続，情報処理学会研究報告, 2017-HPC-160(2017).
[6]Kobayashi, R., Oobata, Y., Fujita, N., Yamaguchi, Y.and Boku, T.: OpenCL-ready High Speed FPGA Net-work for Reconﬁgurable High Performance Computing,Proceedings of the International Conference on HighPerformance Computing in Asia-Paciﬁc Region, HPCAsia 2018, New York, NY, USA, ACM, pp. 192–201(online), DOI: 10.1145/3149457.3149479 (2018).
[7]RIKEN: Himeno Benchmark, (online), available from⟨http://accc.riken.jp/en/supercom/documents/himenobmt/⟩.
[8]Sano, K., Kono, F., Nakasato, N., Vazhenin, A. andSedukhin, S.: Stream Computation of Shallow WaterEquation Solver for FPGA-based 1D Tsunami Simula-tion, SIGARCH Comput. Archit. News, Vol. 43, No. 4,pp. 82–87 (online),DOI: 10.1145/2927964.2927979(2016).
[9]Idomura, Y., Nakata, M., Yamada, S., Machida,M.,Imamura,T.,Watanabe,T.,Nunami,M.,Inoue,H.,Tsutsumi,S.,Miyoshi,I. and Shida,N.: Communication-overlap techniques for improvedstrong scaling of gyrokinetic Eulerian code beyond100k cores on the K-computer,The InternationalJournal of High Performance Computing Applica-tions,Vol. 28,No. 1,pp. 73–86 (online),DOI:10.1177/1094342013490973 (2014).
[10]藤田典久，小林諒平，山口佳樹，朴泰祐，吉川耕司，安部牧人，梅村雅之： 並列 FPGA システムにおける OpenCLを用いた宇宙輻射輸送コードの演算加速，情報処理学会研究報告, 2018-HPC-165 (2018).
c⃝ 2018 Information Processing Society of Japan9
Vol.2018-HPC-167 No.9
2018/12/17