« 超リアルタイムHMD試作に向けた技術検証(中間報告?) | トップページ | FPGAのLUTで作るバイナリニューラルネットのアーキテクチャ考察 »

2018年8月 4日 (土)

FPGAアーキテクチャ上でのバイナリDNNについての考察

 ご承知の通り、世の中ではDL(Deep Learning)の一定の成功の元に、DNN(Deep Neural Network)が大流行です。

 最近ではやや(ビジネス的な)限界点も見え初めて少し落ち着き気味(最初が盛り上がりすぎなだけ?)ではありますが、私としましても無視は出来ない技術ですので、遅ればせながらキャッチアップを試みているところです。
 とはいえ、本当に遅ればせで今更感も強いので、なんとか人がやってなさそうなところを攻めてみようということで、FPGA的な6入力LUTだけを使ったDNNを考察してみたいと思います(Learningは一旦置いておいて、Prediction専用回路を検討しますのであえてDLと言わずにDNNと言っておきます)。
 一般の多値のDNNに関しては、各ノード(ニューロン)に接続された入力に係数を乗じて足しこむという演算が大半を占めますので、巨大な係数テーブルの読み込みと乗算の能力がプロセッサとしての性能を決めてきます。
 したがって、乗算器アレイをぶん回すのがメリットが大きく、GPUが汎用性の高さに対して非常にアーキがマッチしていて性能が出ますし、FPGAでがんばればもっと高効率になるし、専用LSIを作れば超高効率になるわけです。
 一方で、単純に計算機アーキテクチャの側面だけから見ると
  • 乗算はトランジスタ数を食う(一般に乗算器の再利用が必要)
  • 各層の全接続に伴い、係数テーブルが非常に大きい(一般に外部メモリが必要)
  • 演算器の再利用でスイッチなどの補助機構にトランジスタを食われる
  • 演算器再利用でデータの退避復帰を伴うのでリアルタイム性が出し難い
 などの課題も見えてきます。
 そういった中でこれらの問題に対し有力そうなアプローチとして、BinaryNetXNOR-Net などのバイナリ系のネットワークの存在があるようです(英語苦手だし、まじめに論文は読んでないのですが、こちらで概要を知りました)。
 バイナリ化によって分類精度自体は下がるようなので、メインストリーム(如何に精度を上げるために計算機に難しいことをさせるか)からは外れるジャンルなのかもしれませんが、効率やリアルタイム性などが重要になるアプリに関してはメリットになる可能性も大いになるかと思います。

 で、ここまでは既存技術の調査なので、FPGA派の私としては、もう一歩踏み込んで、「いっそ6入力LUTにそのまま適用できるDNNは作れないのか?」という疑問を元に少しアプローチしてみました。
 手元にXilinx社のXC7Z020というFPGAがありますが、DSP(乗算器)数220個に対して、LUTは53,200個搭載している上にDSPより高速です。なので、LUTでネットが組めれば、そもそも係数をLUT内のテーブルに係数をROM的に内蔵して、各ノードごとに専用にLUTを割り当てれば演算器の再利用も不要なのではと、思ってみたわけです(従来のDNNだと外部メモリにデータや係数を置いて、DSPに休み無くそれを供給するための回路を作るのにLUTを使うわけですが)。
 今回の試行のベースにしたのがO'REILLYの「ゼロから作るDeep Learning」本です。こちらソースはgithubに公開されているのですが、第3章に4層のネットでのMNISTに対する学習済みのパラメータがあり、predictionのみを実行するコードがありますので、まずはこのパラメータを借用して初期実験をしてみました(先日、勉強がてらC++でリライトしたり、Eigen使ってみたりしたのででそこから改造しています)
 まず、オリジナルの認識率は 0.9352 です。
 これに対してバイナリ化を試みます。
 まず入力画像が 0.0~1.0 に正規化されていますが、これを 0.0 と 1.0 の2種類に2値化します。
 また出力層の sigmoid や softmax はすべて結果を -1.0 or +1.0 に2値化する演算に置き換えます。
 これで、各ニューロ素子の入出力はすべて2値化されたことになります。
 ただし、このままでは、出力層の10個の中から最大値を探すというようなarg max演算が出来ません。
 そこで一工夫します。
 入力の画像の2値化について、単純
2値化ではなく乱数(std::mt19937)を用いて、2値化を行い、同じものを同じネットに64回(暫定)投げ込みます(当然PCだと計算時間が64倍.....)。
 乱数ですので、同じソース画像から、中間階調部分にディザ的な効果がかかって、
微妙に異なる2値画像が64枚出来るわけですが、このそれぞれを同じネットに入れると結果が微妙に異なるのでそれを集計して、もっとも+1.0の出力回数が多かった素子を採用(要するに多数決)します。
(もともとMNISTが2値に近い画像なので、少しスケーリングしてからディザを考えていたのですが、結果としては何もせずにディザが最善でした
 この作業を行ったときの認識率が 0.5018 でした。もっと悲惨なことになるかと思いましたが、意外にそれっぽい結果が出ています。
 ここでさらに、層間を全結線するのではなく、重みの上位6個以外切り離すということをしてみます。これでFPGAのLUTに嵌る形になるはずです。
 操作としては単純に重さの絶対値でソートして、上位6個以外の重みを0にします。
 で、この時の認識率が 0.3975 でした。
 元データの各層のノード数が 784-50-100-10 ですので、6入力にするとそもそも1層目などは、どのノードからも参照されていないノードが発生しているわけですが、それでも40%近いスコアで残るとは思いませんでした、箸にも棒にもかからないってことはなさそうです(コードバグってないか心配ですが)。
 ひとまず、従来のDLの手法(誤差逆伝播など)で学習した結果を、2値化して利用するという手はそこそこ機能しそうだと言うことは見えてきました。
 で、なんとなく光が見えてきたので、さらに踏み込みたいと思います。
 LUTは万能なロジックなので、単に重みをつけて総和をとる以上の演算が可能です。だったらLUTが可能な演算は何でも許す方向で学習できないか?というところに興味が行きます。実際6入力に対して2^64=18446744073709551616通りの出力が出来るLUT素子は実はかなり高密度な情報が詰め込める気がしています。
で、やってみたのですが、結論から言うと今のところ思っていたよりは意外に悪くない結果がでている気がしています。
 まず、この考え方の最大の問題は学習方法です。
 そもそも演算内容を規定していないので、誤差逆伝播的な数学的手法の適用を思いつけません。
 ということで、まずはベタな手法として、結果を見ながらLUTの内容を確率的に変更していくという方法をとります。
 今回は感触が見たいので、普通に6入力LUTとして動作する64個のテーブルを持ったC++のクラスを作り、これを結合してシミュレーションすると言う計算速度的に早くなる余地の無い馬鹿コードで試します。
 入力はMNISTを乱数で2値化した28x28=784bitの入力を前提として、784-200-50-10 の4層で組みました。なのでLUTとしては 200+50+10 の 260個というFPGA化したなら微小回路になる構成です(まあ、PC計算だと無駄だらけで重いんですが)。
 学習のやり方として、
  • MNISTの train と test で学習用と評価用データセットは分ける
  • LUTの中身は乱数で初期化(この段階で当然正解率はほぼ10%)
  • LUTの結線も一旦乱数で前の層の6個にランダムに接続
  • 特定のLUTに着目して、学習データのサブセットを流し、LUTの入力に対して0出力の場合と1出力の場合で結果がどうなるか統計を取る
  • 統計結果に伴って、乱数を使って、結果の改善が良い方に高確率で書き換わるように学習を実施
  • この操作を対象LUTを変えながらひたすら繰り返す
  • 定期的に評価用データセットで状況を見る
といった、力技の実装。
で、流してみたところ、正解率は上下しながら徐々に上昇していき、計算した範囲ではmaxで、0.4749を記録しました。
 デジタルDNNの視点でLUTを再度考えてみると、入力数が6に制限されていることを除けば、デジタルとしては万能ニューロンなわけですから、これを回路構成のままニューラルネットワークとして機能させることが出来れば、実装効率の観点では面白いものになる可能性がないかと思っている次第です。
 
 もう少し深ぼりしたいところですが、今のプログラムだと、FPGAなら1サイクル数万並列でやれる計算を、「6回ランダムアクセスでロードしてきてテーブル引いて出力」というコードのfor文回している状態なので、死ぬほど非効率です。
 ネットの構成を変えたり学習のやり方を工夫したりいろいろ試行錯誤するには今の環境だと重たすぎてやる気が起きないので、ここは改善が必要そうです。
 GPUもあまり良いもの持ってないので、一旦扱いやすい OpenMP + Intel-SIMD(AVX2) とかで書き直してから次のステップに進もうかと考えている次第です。
 FPGAだと自由にルーティングできますが、プロセッサだと6個のランダムアクセスになると言うところが意外に性能向上の罠でして、_mm256_i32gather_epi32などの命令を初めて使うことになりそうです。
 
 先は長そうですが、しばらく遊べそうなネタではありそうです。さて、どうなることやら。
<追記>
 今、0.6186 とかのスコアがでとります。いろいろバグ抱いていそうなので、まだまだ改善するかもです。
<さらに追記>
入力層784bitに対して  200-100-50-10 の LUT:360個の回路で学習中で、現在0.7221の正解率。いろいろ学習時の係数で収束速度が大きく違う模様。
 配置配線したときに配線が課題にならないかどこかで試しておきたい。
<さらに追記[2018/08/05]>
400-300-200-50 のNET で 0.7914 まできました(一晩ほど計算)。
LUT増やせばある程度性能は上げられそう。オリジナルの乗算器型と比べるとロジック規模は格段に小さいはずなので、落としどころが重要ですね。CNN的な構成も検討候補になってくると思います。

« 超リアルタイムHMD試作に向けた技術検証(中間報告?) | トップページ | FPGAのLUTで作るバイナリニューラルネットのアーキテクチャ考察 »

FPGA」カテゴリの記事

Deep Learning」カテゴリの記事

コメント

この記事へのコメントは終了しました。

トラックバック


この記事へのトラックバック一覧です: FPGAアーキテクチャ上でのバイナリDNNについての考察:

« 超リアルタイムHMD試作に向けた技術検証(中間報告?) | トップページ | FPGAのLUTで作るバイナリニューラルネットのアーキテクチャ考察 »