OpenCL meets FPGA #2 最適化編 #FPGA

OpenCL meets FPGA #1 入門編では、OpenCLを用いたFPGAプログラミングについて紹介しました。

FPGA向けOpenCL実装の紹介
Altera SDK for OpenCL (AOCL) ツールチェインの使い方
AOCLにおけるOpenCL C言語とFPGAデザインの対応

本記事では、引き続きAOCLを利用して、OpenCL C言語を用いてFPGA上に高性能なOpenCLカーネルパイプラインを構築する方法について、実例を交えながら説明します。

潜在的ボトルネック

最適化の本質は、ボトルネックを取り除くことです。性能向上のためのテクニックは多種多様ですが、本質は変わりません。では、FPGAにおけるボトルネックとは何でしょうか？

その問いに答えるには、最適化対象のハードウェアについて知る必要があります。例として、Nallatech社製のPCIe385というOpenCL対応FPGAボードの構成を見てみましょう。

出典：Nallatech PCIe385 データシート

ボード上には、FPGAとしてAltera StratixV、メインメモリはDDR3 SDRAMが実装されています。このボードはホストマシンのPCI Expressスロットに接続し、OpenCLデバイスとして使用できます。

OpenCL APIを用いて、このボード上で計算を行う手順はというと：

clEnqueueWriteBufferでデバイス側のメモリにデータを転送
clEnqueueNDRangeKernelでカーネル実行
clEnqueueReadBufferでホスト側のメモリにデータを転送

基本はこれだけです。では、それぞれの操作の中で何が行われているか？絵にすると、こうなります。

どこが遅そうか、わかりますか？

1, 3はPCI Express経由のデータ転送です。PCIe385の対応バス規格はPCIe3.0x8で、ホストの性能にもよりますが、実効帯域はだいたい2~3GB/sくらいです。

2のカーネル実行は、大きく分けて：

2a. グローバルメモリアクセス
2b. カーネルパイプライン実行

この2つの要素で構成されています。OpenCL Cの規格に準拠するなら、OpenCLカーネルのI/Oはグローバルメモリを経由せざるを得ません。PCIe385の標準BSPでは、グローバルメモリはボード上のDRAMに配置されます。このメモリとStratix Vは2チャネルのDDR3-1600で繋がれているので、理論値は25.6GB/sです。

I/O周りの理論性能を、同じような立ち位置にあるディスクリートGPUとくらべてみましょう。

	PCI Express 帯域 (GB/s)	DRAM 帯域 (GB/s)
NVIDIA Tesla K40	16.0	288
AMD FirePro S9150	16.0	320
Nallatech PCIe385	8.0	25.6

ホスト経由で単なるアクセラレータとしてFPGAボードを使う場合、I/O性能に律速するような単純なタスクのスループット勝負では、GPUに対してかなり分が悪い事がわかります。
電力性能比にすればトントンですが、コスト性能比が悪すぎます。FPGAも数が出れば安くなるでしょうが、コンシューマ向けに大量にチップをばらまいているGPUに正面切って物量で挑むのは、現時点では茨の道と言わざるを得ません。

では、ディスクリートFPGAが生きる道はどこか？率直にいうと：

deeeeeepなパイプラインが生きる、Compute Intensiveな問題
外部I/Oと直結してレイテンシ勝負

メモリI/Oネックになった瞬間に負けるので、それ以外の場所で勝負になります。1はまあ、順当な感じですが、2が謎ですね。外部I/Oって、何でしょうか。

ここまで、意図的に無視してきましたが、PCIe385にはSFP+という光インタフェースが実装されています。

SFP+の物理層は10Gbpsで通信できて、385にはこれが2つ搭載されています。

ここで、お前はさっき自分で書いたことも忘れたのかというツッコミが飛んできそうですが・・・

OpenCL Cの規格に準拠するなら、OpenCLカーネルのI/Oはグローバルメモリを経由せざるを得ません。

そう、OpenCLカーネルとつなぐには、そもそもOpenCLの枠組みからいくらか飛び出す必要があります。またぞろおかしなベンダ拡張か、ポータビリティとか夢物語じゃねーか、という声が聞こえてきそうですが、理想と現実は常に背中合わせであり、我々開発者はそれに向き合ってこそ、前に進むことができるのです！

ということで、今回は詳しく触れませんが、AlteraのOpenCL拡張を使うと、SFP+のような外部I/Oを使うことができます。ここと直結できれば、バスやメモリ、OSなどが介在しない世界で、OpenCLカーネルパイプラインにデータを流しこむことができ、FPGAの真の力（！）を開放することができます。

えーと、よくわからなくなってきましたが、今日は最適化の話をしたいので無理やりまとめると：

以上の問題を把握した上で、カーネルパイプラインの最適化を始めるべき

え？まとまってない？聞こえない・・・

パフォーマンス・メトリクス

カーネルパイプラインを最適化するには、

スループット
レイテンシ
リソース
周波数

これらがどう関係しているか、理解しておくことが重要です。スループットとレイテンシに関しては、前回記事を見て下さい。

リソースは、FPGA上で組み合わせ可能なプリミティブをどの程度使用すればカーネルパイプラインが構成できるか、という目安になります。プリミティブにはロジックエレメント、レジスタ、ブロックメモリ、DSP、などがあり、ターゲットのFPGAチップによってその構成は異なります。

リソースはカーネルコードが複雑で大きくなるほど消費量が上がるので、闇雲にコードを書くとあっという間にコンパイルできない（あるいはPlace&Routeに超〜時間がかかる）OpenCLカーネルが出来上がります。だいたい余裕をもって、90%くらいに抑えるようにしましょう。AOCLはVerilogに変換した段階で大体のリソース使用量を見積もることができて、

aoc -c --report vecadd.cl

こうすると、

+--------------------------------------------------------------------+
; Estimated Resource Usage Summary                                   ;
+----------------------------------------+---------------------------+
; Resource                               + Usage                     ;
+----------------------------------------+---------------------------+
; Logic utilization                      ;   22%                     ;
; Dedicated logic registers              ;    8%                     ;
; Memory blocks                          ;   23%                     ;
; DSP blocks                             ;    0%                     ;
+----------------------------------------+---------------------------;

こんな感じのレポートを出してくれます。これはあくまでフェーズ1における見積で、ビットストリームまで行くとリソース消費は下がることが多いです。最終結果はコンパイルディレクトリ以下のacl_quartus_report.txtを見て下さい。

周波数は、最終的に動作できるカーネルパイプラインの周波数で、やはり回路の複雑さに依存して下がっていきます。なので、頑張ってスループット上げたのに回路複雑になって周波数下がってトータル性能は下、みたいなことがよくあります。経験的には、周波数の上限はStratixV系だと大体250MHzくらいです。

逆に、そんなに性能いらなくて消費電力下げたい（測ったことないのでどれくらい下がるのかわかりませんが）、みたいな時には上限をキャップすることはできます。

aoc --fmax 100 vecadd.cl

関係性をまとめてみました。

性能を測定する

正確な時間計測なく性能を語ってはいけません。（とくにC++03以前では）ポータブルで高精度な計測をするのは案外面倒だったりするのですが、幸いOpenCLにはカーネルの実行時間を取得するAPIが備わっています。

prof.cc

// プロファイル取れるキューを作る
cl::CommandQueue queue(context, devices[0], CL_QUEUE_PROFILING_ENABLE);


// 測定対象カーネルの起動時にイベントを渡す
cl::Event event;
queue.enqueueNDRangeKernel(kernel,
                           cl::NullRange,
                           cl::NDRange(size),
                           cl::NullRange,
                           NULL,
                           &event);
event.wait();

// カーネル実行時間をusで表示
cl_ulong start = event.getProfilingInfo<CL_PROFILING_COMMAND_START>();
cl_ulong end = event.getProfilingInfo<CL_PROFILING_COMMAND_END>();
std::cout << static_cast<double>(end-start)*1e-3f << " us" << std::endl;

OpenCLのプロファイリング・キューによる測定結果はns単位で、デバイス側のタイムスタンプが取れます。

理論性能を求める

カーネルパイプラインは、その構成が決定した時点で理論性能も求まります。CPUと違って、OpenCL界のFPGAがパイプラインストールする要因は、メモリアクセスによるバス混雑と、データ依存によるLSUのストールしかありません。また、カーネルパイプラインを構成した時点で、最大で必要な帯域も自明です。なぜなら、そのパイプラインを1サイクル1ワークアイテムで動かすには、（つまりパイプライン中の全LSUが同時に動作するには）どれくらいのメモリアクセスが必要かは、静的に解析できるからです。

ループの無いカーネルパイプラインの理論性能の計算は以下のとおりです。

全ワークアイテム数 / 動作周波数(Hz) = 実行時間（sec）

単純ですね。最適化するときは、常にこれを計算して、実測値との乖離が無いか確認しておきましょう。

次は、いよいよ、カーネルパイプラインの性能を上げる方法に入ります。

スループットを上げる

スループットは、単位時間あたりに処理できるデータ数です。パイプラインの深さに対してワークアイテム数が十分に多い場合、結局はスループットを上げないと性能は上がりません。スループットを上げる、ということは、1ワークアイテムあたりが処理できるデータ数を増やすか、サイクルあたりのワークアイテム数を増やすか、どちらかです。これには3つ方法があって：

カーネルパイプラインを多重化する
カーネルパイプラインが処理できるデータ数を増やす
ループアンロールする

順に見て行きましょう。

カーネルパイプラインを多重化する

これはOpenCLカーネルにnum_compute_unitsアトリビュートを指定することで達成できます。

vecadd.cl

__kernel
__attribute__((num_compute_units(2)))
void vecadd(__global const int *a, __global const int *b, __global int *c)
{
  const int i = get_global_id(0);
  c[i] = a[i] + b[i];
}

上の場合、カーネルパイプラインが2つ構成されて、ワークグループ単位で振り分けられます。当然、カーネルパイプライン部分が消費するリソースは2倍になりますが、スループットも2倍です。ワークグループ単位で振り分けが行われるので、複数のワークグループができるようにカーネルを起動しましょう。

カーネルパイプラインが処理できるデータ数を増やす

結論から言うとベクトル化しましょう。まず、アトリビュートを使う方法。

vecadd.cl

__kernel
__attribute__((num_simd_work_items(2)))
__attribute__((reqd_work_group_size(64, 1, 1)))
void vecadd(__global const int * restrict a, __global const int * restrict b, __global int * restrict c)
{
  const int i = get_global_id(0);
  c[i] = a[i] + b[i];
}

num_simd_work_itemsには2, 4, 8, 16のいずれか、ベクトル化数を指定できます。ベクトル化を自動でやるには、reqd_work_group_sizeにワークグループのサイズを指定して下さい。ワークグループのサイズが可変だと、色々静的に決まらないのでベクトル化できません。特に可変にしたい理由がなければ常に指定しておいて下さい。

また、restrict修飾子も大事です。上の例なら、a, b, c由来のアドレスは同じ領域にならないということをコンパイラに教えています。バッファがエイリアスしなくてコンパイラの最適化を期待したいなら、restrict修飾子は常につけるべきです。

コンパイラは色々解析をしてベクトル化しようと頑張りますが、当然、静的解析の限界とかそういうので、頑張れない場合もあります。また、自動ベクトル化の結果、パイプラインが美しくないということもありえます。そんな時は・・・手動でベクトル化しましょう。

vecadd.cl

__kernel
void vecadd(__global const int2 * restrict a, __global const int2 * restrict b, __global int2 * restrict c)
{
  const int i = get_global_id(0);
  c[i] = a[i] + b[i];
}

これくらいなら簡単ですね。ホスト側のワークアイテム数をちゃんと調整するのを忘れないようにしましょう。

また、静的解析できそうな単純なループで書くとコンパイラがベクトル化してくれることもあります。

vecadd.cl

__kernel
void vecadd(__global const int * restrict a, __global const int * restrict b, __global int * restrict c)
{
  const int i = get_global_id(0);
  #pragma unroll
  for (int j=0; j<2; ++j) {
    c[i*2+j] = a[i*2+j] + b[i*2+j];
  }
}

これくらいならコンパイラが勝手にやりますが、pragma unrollでループアンロールしています。ループのほうが汎用的に記述できますが、中間コード見るなりして思った通りベクトル化されているか確認したほうが良いでしょう。.bcをポイと渡されて読める人には、こっちの方がお勧めです。

パイプラインをまるっと並べる1の方法に比べて、ベクトル化は消費リソースが少なくてすむというメリットがあります。基本はベクトル化を試して、どうしてもできない時はパイプラインの多重化を試しましょう。

ループアンロールする

前回、ループは循環のあるパイプラインに変換されて、スループットが下がるという話をしました。リソースが許す限り、アンロールするのが良いでしょう。

loop.cl

__kernel
void loop(__global const int * restrict src, const int num, __global int * restrict dst)
{
  const uint i = get_global_id(0);
  int sum = 0;
  #pragma unroll 2
  for (uint j=0; j<num; ++j) {
    sum += src[num*i+j];
  }
  dst[i] = sum;
}

pragma unrollにはアンロール数を指定できて、上の例では、ループ回数は半分で済むので、スループットは2倍です。もちろん、ループブロック部分の演算リソースの消費量も2倍なので、空きリソースとのせめぎ合いになります。

ちなみに、何もつけないと全部アンロールしようとしますが、ループ回数が解析できない場合アンロールは行われません。

loop.cl

__kernel
void loop(__global const int * restrict src, const int num, __global int * restrict dst)
{
  const uint i = get_global_id(0);
  int sum = 0;
  #pragma unroll
  for (uint j=0; j<num; ++j) {
    sum += src[num*i+j];
  }
  dst[i] = sum;
}

Compiler Warning: Full unrolling of the loop is requested but the loop bounds cannot be determined. The loop is not unrolled in kernel loop

いい感じに使い分けましょう。

レイテンシを下げる

レイテンシ（＝パイプラインの深さ）を下げたい理由は2つあって：

レイテンシ下げることで消費リソースを減らして、そのぶんをスループット上げるのに使いたい
そもそもレイテンシが性能に直結している

2の場合というのは、パイプラインの深さに対してワークアイテム数が少ない時か、外部I/Oと直接繋いでレイテンシが超大事という時ですね。では、どういう手段があるか。基本方針は

コードをシンプルにする

とすごい単純なんですが、これでは身も蓋もないのでAOC特有のテクニックを幾つか紹介します。

浮動小数点数の融合演算を許す

みなさんが普段何気なく使っているfloatやdoubleといった浮動小数点数は、IEEE754という規格によって、丸め方向やNaN/Infなどの特殊な値の取り扱いなど、厳密にその挙動が定められています。OpenCLの浮動小数点数はこれに従うことが規格で決められていて、AOCもこれに則った変換を行います。

AlteraのFPGAは可変精度乗算器をハードマクロとして搭載していて、乗算じたいはLEを使わないでできます。ただ、IEEE754に準拠するためには、丸め等の後処理が入ります。これが1回ならよいのですが、演算ツリーが大きくなってくると演算回数に比例してレイテンシが上がってしまいます。そこで、複数の浮動小数点数演算が連なっている場合に、途中のIEEE754準拠のための色々な処理を省略する、というオプションがあります。

aoc --fpc vecadd.cl

ドキュメントには、丸め等サボる代わりに仮数部を1ビット増やしてあるので真値に対する精度はこっちのほうが上だよー、とか書いてあります。ただしIEEE754に準拠しない上、ポーティング元のコードと厳密一致が求められる場合には使えません。

まあ、1ulp間違ったら駄目な計算にそもそも浮動小数点数なんて使わないから問題ない、ですよね？

浮動小数点数演算のリオーダーを許す

浮動小数点数演算が、演算順序によって結果が変わることを知らないで許されるのは小学生までですが、レイテンシを重視する場合にはこれがネックになることがあります。

これらのコードは：

add-unbalanced.cl

((a * b) * c) * d

add-balanced.cl

(a * b) * (c * d)

それぞれこんな雰囲気の回路になります。

バランスしている方がレイテンシが少なくてすみます。ということで、add-unbalancedをadd-balancedに自動的に変換する事を許すオプションが、--fp-relaxedです。

aoc --fp-relaxed vecadd.cl

ツリーに引き算が入っていると、変換によってお互いに近い値を引き算して思わぬ誤差を生む、みたいなこともあるので、気をつけて使いましょう。

使ってはいけない命令たち

div（おもい）
atomic（おそい）
async（意味が無い）
fp math intrinsic（リソース爆発）

頑張って避けましょう。

まとめ

今回紹介した最適化方法は全て公式ドキュメントに記載されています。詳しく知りたい方は

あたりを読みましょう。

今回説明しなかった内容に、

メモリアクセスまわりの最適化
Altera Channel拡張を使用した最適化

があります。機会があれば書くかもしれません。

最後に・・・

プロセス微細化の限界、ムーアの法則、ダークシリコン問題、etc・・・よってCPUの時代はもう終わり、FPGAが銀の弾丸である！といった論調を見かけますが、FPGAも同じ半導体テクノロジで実装されたチップに過ぎません。3年前にGPUでも繰り返された歴史です。まずは、w_o先生による、Debunking the 100X GPU vs. CPU Myth (訳:GPUなら100倍速いという神話を覆すを正座して読みましょう。

その上で、盲信して賞賛するのではなく、特性を知り適切なアプリケーションを適切な方法で実装したならば、FPGAはあなたのアプリケーションを加速する武器の1つになりうるでしょう。