TadaoYamaokaの日記

山岡忠夫Homeで公開しているプログラムの開発ネタを中心に書いていきます。

PUCTの定数のベイズ最適化

AlphaZeroの論文では、PUCTの定数C_{PUCT}を以下の式で、親ノードの訪問回数に応じて動的に調整を行っている。

U(s,a)=C(s)P(s,a)\frac{\sqrt{N(s)}}{1+N(s,a)} \\
C(s)=\log{\frac{1+N(s) + C_{base}}{C_{base}}}+C_{init}
この式で現れる定数C_{base}C_{init}は、疑似コードでは以下のように定義されている。

    # UCB formula
    self.pb_c_base = 19652
    self.pb_c_init = 1.25

私が実験しているdlshogiでも上記の式を適用することで効果があるか試してみた。

報酬によるスケーリング

AlphaZeroは、報酬(負けと勝ち)を-1、1で与えているが、dlshogiでは、0、1で与えているため、スケールを変更する必要がある。
定数C_{PUCT}の理論値は、報酬が[-1,1]の場合は2で、[0,1]の場合は1となる。
したがって、およそ半分になるように調整すればよい。

測定結果

以下の通り、手動でいくつかの組み合わせで、GPSfishと1手3秒で100回対局した勝率を測定した。

c_base c_init 勝率
変更前(C_PUCT=1) 54%
25000 0.8 65%
20000 0.7 62%
30000 0.85 56%
20000 0.8 57%
30000 0.8 66%

考察

いずれの組み合わせでも、変更前のC_{PUCT}が固定値1の場合より、勝率が高くなっている。
勝率が最大の組み合わせは、c_base=30000、c_init=0.8であった。
C(s)をグラフにすると以下のようになる。
f:id:TadaoYamaoka:20181213001907p:plain

自己対局でモデルの学習を長時間行っても勝率は少しずつしか変わらなかったが、探索のパラメータを変えるだけで10%以上勝率が上がることがわかった。
今までは探索パラメータを手動で適当に決めていたが、強くするには探索のパラメータの調整も重要であると気づかされた。

そこで、ベイズ最適化を使って、探索パラメータをちゃんと調整することにした。

ベイズ最適化

ベイズ最適化を使うことで、既知の説明変数と目的変数から、確率的に目的変数が最大となる説明変数の組み合わせ候補を求めることができる。
ベイズ最適化の理論を正確に理解して実装するのは大変なので、こちらのページにあったPythonのコードを使用させてもらった。

以下のようにして、次の実験候補を求めることができる。

import numpy as np

from bayesianoptimization import bayesianoptimization

X = np.random.rand(100, 2) * [1, 3] + [0.5, 1]
X_train = np.array([[0.8, 2.5], [0.7, 2.0], [0.85, 3.0], [0.8, 2.0], [0.8, 3.0]])
y_train = np.array([0.65, 0.62, 0.56, 0.57, 0.66])

selected_candidate_number, selected_X_candidate, cumulative_variance = bayesianoptimization(X_train, y_train, X, 2)

selected_X_candidateに説明変数の候補が格納される。

array([0.70598003, 2.04039803])

候補となった条件で、実験を行い、その結果を追加して、同様に次の候補を求めて実験すればよい。

変数の正規化

c_baseの値の桁がc_initに比べて大きいため、適切に正規化しておく必要がある。
上記のコードでは、c_baseを1/10000にしている。勝率も%ではなく、[0,1]の範囲で与えている。
Xは、ランダムにサンプリングした仮想的な候補だが、サンプリングの範囲は、c_baseを[1.0,4.0]、c_initを[0.5,1.5]の範囲でサンプリングしている。

ベイズ最適化で求めた候補の実験結果

実験した結果は以下の通りとなった。

c_base c_init 勝率
20403.9803 0.70598003 59%

予測はしていたが、1回の実験では最大にはならなかった。
そもそも100回の対局では誤差も大きく、確率的な事象であるため、1回で最適な値は求まらない。

繰り返し実験を行う必要があるが、今後は自動的に調整できるような仕組みを構築したい。

AlphaZeroの論文

Science誌に掲載された論文は、新しい対局条件での結果と棋譜の公開がメインで技術的な内容は、昨年のarXivで公開された論文とほとんど差分はありませんでした。
DeepMindのページのリンクからダウンロードできるOpen Access versionのMethodsでは、技術的な内容が追加されており、興味深い内容も記載されていました。

Prior Work on Computer Chess and Shogi

TD Gammonwasについて追加されています。

Domain Knowledge

引き分けの条件がチェスと将棋では512手、囲碁では722手であることが追加されています。
定跡と終盤データベースを使用していないことが追加されています。

Search

MCTSアルゴリズムの説明が追加されています。
以前のAlphaGoからPUCTの定数が変更されており、親ノードの訪問回数応じた関数になっていました。

以前

U(s,a)=C_{PUCT}P(s,a)\frac{\sqrt{N(s)}}{1+N(s,a)}

今回

U(s,a)=C(s)P(s,a)\frac{\sqrt{N(s)}}{1+N(s,a)} \\
C(s)=\log{\frac{1+N(s) + C_{base}}{C_{base}}}+C_{init}

C(s)は探索の割合(exploration rate)で、探索の時間とともにゆっくり増加し、探索が高速な場合は、本質的には一定であると説明されています。

定数の値は、公開された疑似コードをみれば分かります。(Additional FilesのData S1に含まれるpseudocode.py)

    # UCB formula
    self.pb_c_base = 19652
    self.pb_c_init = 1.25

と定義されています。

グラフにすると以下のようになります。
f:id:TadaoYamaoka:20181208181048p:plain

訪問回数が少ない内は探索結果を重視して、訪問回数が増えるとボーナス項の重みが大きくなり、未知の手が選択されやすくなると解釈できます。

Representation

ポリシーの予測した手について、合法でない手は確率を0にして、残りの合法手で再正規化するという説明が追加されています。

Architecture

ネットワーク構成の説明が追加されています。

  • AlphaGo Zeroでは、ポリシーの出力に全結合層が使われていましたが、チェスと将棋では、畳み込み層であるという説明が追加(ポリシーの出力層の畳み込みのカーネルサイズについては記載されていません。)

チェスと将棋でポリシーの出力が畳み込み層になっていることから、AlphaGo Zeroでポリシーの出力が全結合層になっていた理由は、パスを表現するためであったと解釈できます。(個人的に知りたかった内容でした。)

バリューの出力層は、1フィルター、カーネルサイズ1×1の畳み込み層の後に、256ユニットの全結合層と1ユニットの全結合層で、tanhで出力しています。
チェスと将棋でも、畳み込み層は、1フィルターになっていました。
駒の種類が多いため、フィルタ数は多くしていると予測していましたが、そうではなかったようです。
バリューに関してはフィルターを1つにして、重みを共有した方がよいのかもしれません。

Configuration

学習率について詳細が追加されています。
学習率を0.2から0.02、0.002、0.0002に落とす条件が、チェスや将棋の場合は100、300、500,000ステップ後で、囲碁の場合は0、300、500,000ステップ後という条件が追記されています。

Opponents

  • 対局相手のソフトの条件の記載が追加

定跡の仕様有無について一部で疑問が上がっていましたが、将棋は、やねうら王の標準定跡を使用していることが記載されています。
ハードウェアなどの詳細な条件が記載されていましたが割愛します。

Match conditions

対局条件について内容が追加されています。
初期局面から対局を行い、将棋の場合は2017年WCSCと同じ持ち時間10分、1手10秒加算の条件です。
AlphaZeroの時間制御は、単純に残り時間の1/20のようです。
ポンダーは、思考時間の評価を行うため、双方未使用の設定です。
投了の条件は、elmoは-4500が10回連続の場合、AlphaZeroはバリューが-0.9になった場合です。
1,000試合で勝率を評価しています。

Elo ratings

Eloレーティングの測定方法について内容が追加されています。
1手1秒の対局結果から評価したことが書かれています。

感想

ポリシーの出力層の構成が知りたかったことの一つだったので、有益な情報が得られたと思います。
PUCTの定数が訪問回数による関数に代わっていたことは、dlshogiでもどれくらい効果があるか試そうと思います。
疑似コードがPythonで実装されていたので、AlphaZeroもPythonで実装されているのか気になりました。
まだ全部読めていませんが、疑似コードを読むと自己対局から訓練の流れや、データの格納方法とかが推測できて楽しめます。

2018/12/9 追記

疑似コードを見ると、学習のハイパーパラメータがわかります。
自己対局して学習に使う局面数は、700,000ステップ(1ステップは4,096局面)と記載されていましたが、過去何局からサンプリングしているかは記載されていませんでした。(AlphaGo Zeroの論文には、most recent 500,000 gamesと記載されていました。)
疑似コードには、

    ### Training
    self.training_steps = int(700e3)
    self.checkpoint_interval = int(1e3)
    self.window_size = int(1e6)
    self.batch_size = 4096

と定義されており、ReplayBuffer.save_gameで、保存するゲームの数の上限がwindow_sizeになっていることから、直近1,000,000局からサンプリングしていることがわかります。

タスクマネージャーのメモリ使用量

将棋AIの学習のためにGPUを増やしたところ、まったく速度が上がらないどころか低下するという事象が起きて、ここ数日原因を調べていました。
メモリスワップが起きていたことが原因で、わかってしまえば単純な話でしたが、Windowsのタスクマネージャーの仕様の理解が不足していたため、数日悩むことになりました。

物理メモリを食いつぶしていたので、Windows 10のタスクマネージャーの詳細タブでプロセスのメモリ使用量を見ると、プロセスを起動してから、メモリ(プライベートワーキングセット)の数字が徐々に上がっていたため、メモリリークを起こしているのではないかと疑いました。

プライベートワーキングセットとは

実は、タスクマネージャー表示されるメモリ(プライベートワーキングセット)は、プロセスに割りあてられたページの物理メモリのサイズを示しており、C++のnewで確保したメモリは、仮想メモリが割り当てられるだけで、実際にアクセスするまではOSに物理メモリが割り当てられないということが今回調べてわかりました。
詳しい人によっては、常識レベルの話で、自分も仮想メモリの動作は知識としては知っていたはずですが、タスクマネージャーの数値の意味を理解していませんでした。

なお、仮想メモリのサイズを見るにはタスクマネージャーで、列の選択で「コミット サイズ」を表示する必要があります。


f:id:TadaoYamaoka:20181121210733p:plain


ということで、徐々に上がる=メモリリークと勘違いしたため、コードを変えたりして原因を探してもリークしている箇所が見つからず数日時間を浪費しましたorz。
しかし、おかげでタスクマネージャーと仮想メモリについて理解が深まったので良いとします。

プログラムの改良

自己対局プログラムは、1つのGPUごとにエージェント複数割り当てて、それぞれのエージェントに別々のハッシュを割り当てていたので、メモリが大量に必要となる構成になっていました。
以前に、エージェントを全て別スレッドで動かしていたときの構成を引きずって、そのような構成になっていましたが、現在はシングルスレッドで直列に動かしているため、ハッシュを共有することで、メモリ使用量を減らすように修正しました。
メモリスワップが起こらないサイズに調整したことで、安定した速度で自己対局が行えるようになりました。

将棋でディープラーニングする その53(価値ネットワークの精度向上)

作成している将棋AIは、現在自己対局による強化学習を続けています。
floodgateの棋譜との一致率は徐々に上がっており、少しずつですが強くなっています。
48サイクル回したところで、GeForce 1080を1枚搭載したノートPCで、GPSFishに1手3秒で勝ち越すようになりました。

対局数100 先手勝ち57(56%) 後手勝ち41(41%) 引き分け2
selfplay048temp100
勝ち58(57%) 先手勝ち33(33%) 後手勝ち25(25%)
GPSfish 0.2.1+r2837 gcc 4.8.1 osl wordsize 32 gcc 4.8.1 64bit
勝ち40(40%) 先手勝ち24(24%) 後手勝ち16(16%)

有意水準以下なのでまだ有意に強くなったとはいえませんが。

自己対局はGPUを2080 Tiに交換したのとFP16化したことで、1サイクル3.5日かかっていたのが2.5日になりました。
それでも時間がかかるので成長速度はゆっくりです。
AlphaZeroのようにわずか2時間でとはいきません。
AlphaZeroが2時間に対して、こちらは1年近くかけています。
しかも、ネットワークの規模はResNet10ブロックなので1/4です。(AlphaZeroは40ブロック)

ということで、もう少し効率よく精度をあげる方法を模索したいと思っています。
囲碁の方では、価値ネットワークにプーリングを使用することが試されているようなので試してみました。

dlshogiでは、方策ネットワークを価値ネットワークを同じネットワークにして、マルチタスク学習を行っているため、価値ネットワークのみに単純にプーリングを適用できない構成になっています。
そこで、出力側の方策ネットワークを価値ネットワークで分岐させている部分にプーリングを適用して実験してみました。

現在のネットワーク構成

方策ネットワークと価値ネットワークで共通部分
  • 10ブロックのResNet(畳み込み2層)
  • フィルターサイズ3×3
  • フィルター枚数192
  • BatchNormalization
  • ドロップアウト
  • プーリングなし
方策ネットワークの出力部分
  • フィルター1×1の畳み込み
  • フィルター枚数2187:駒の種類と移動方向の組み合わせ
  • 位置ごとに異なるバイアス
価値ネットワークの出力部分

1層目

  • フィルター1×1の畳み込み
  • フィルター枚数2187:方策ネットワークと同じ

2層目

  • 全結合層:ノード数256

3層目

  • 全結合層:ノード数1

テストパターン

プーリングの他に、ストライドありでオーバーラップさせた畳み込みや、パディングなしで画像サイズを削減した畳み込みも試した。

test1 現在の構成
test2 1層目をストライド3、フィルターサイズ3×3、フィルター枚数192×3に変更
test3 最後のブロックを方策と価値で分けて、価値の方をストライド3、フィルター5×5、フィルター枚数192×3の畳み込みの後に、フィルター3×3、フィルター枚数192*3の畳み込みに変更。ResNetブロックの後、フィルター枚数192*3の畳み込みとフィルター枚数3×3、パディングなしの畳み込み。その後全結合層。
test4 1層目をストライド3、フィルターサイズ5×5、フィルター枚数192×3に変更し、その後にサイズ3×3の平均プーリング
test5 1層目をストライド3、フィルターサイズ5×5、フィルター枚数192×3に変更
test6 1層目はそのままで、その後に、サイズ3×3の平均プーリング
test7 1層目をストライド3、フィルターサイズ5×5、フィルター枚数192×3×3に変更
test8 1層目はそのままのものと、ストライド3、フィルターサイズ5×5、フィルター枚数192×3の畳み込みの出力を連結(concat)
test9 1層目をストライド3、フィルターサイズ5×5、フィルター枚数はそのままに変更
test10 共通部分のフィルター枚数を256に変更

測定結果

ランダムに初期化されたモデルから、自己対局で生成した1000万局面を毎回シャフルして5回測定した。

パターン 平均 中央値
方策の一致率 価値の一致率 方策の一致率 価値の一致率
test1 0.340667506 0.646848387 0.3411726 0.647756696
test2 0.340076311 0.648765304 0.3398553 0.6486792
test3 0.340311587 0.648855639 0.34040904 0.64919275
test4 0.340737118 0.649633029 0.34051056 0.64909655
test5 0.341156782 0.648948858 0.340723008 0.648917139
test6 0.340287269 0.64832114 0.340290305 0.6485846
test7 0.340536236 0.648031611 0.340501636 0.6472785
test8 0.341074753 0.649948122 0.340924477 0.649788767
test9 0.340526263 0.64853402 0.340724289 0.648599923
test10 0.345498568 0.650716498 0.34547478 0.6507526

それぞれの列で上位3つに色を付けた。

考察

モデル構成そのままでフィルター枚数を増やしたtest10が最も精度が高くなっているが、学習時間が1.5倍近く増えているため、参考値扱いである。

test10を除くと、test8が最も価値の精度が高かった。
test8は、チャネル方向に重み共有した畳み込み(1×1フィルター)と、空間方法に重み共有した畳み込み(オーバーラップさせたストライド)の両方を行って、それぞれの連結している。
現在のモデルは、1×1の畳み込みで、チャネル方向の重み共有で、空間の情報を残している。

test2~test9は何らかの空間方向の重み共有を行っており、test1よりも価値の平均の精度が高くなっている。

このことから、実験回数がそれぞれ5回のみで、収束するまで学習していないので、はっきりは言えないが、価値ネットワークには空間方法の重み共有が有効である可能性がある。
また、test8が最も精度が高かったことから、チャネル方向の重み共有も有効で、併用すると効果が上がる可能性がある。

オーバーラップしながらストライドした畳み込み(test9)と、平均プーリング(test6)では、精度に違いは見られなかった。

まとめ

価値ネットワークは、チャネル方向(1×1の畳み込み)にも空間方向(ストライドあり畳み込みもしくはプーリング)にも重み共有を行った方が精度が上がる可能性がある。


新しいモデルに切り替えて強化学習を試したいが、今からモデルを変更して学習をやり直すのも大変なので、しばらくは同じモデルで強化学習は進めるつもりです。

GeForce RTX 2080 TiでTensorCoreを使う(その2 FP16対応)

前回GeForce RTX 2080 TiのTensorCoreによる推論の速度を、cuDNN 7.2で追加されたCUDNN_TENSOR_OP_MATH_ALLOW_CONVERSIONオプションを使用して、内部的にFP32からFP16に変化させることで測定を行った。

今回は入力から出力までを、FP16にして測定してみた。

FP16対応

cuDNNをFP16で使用する方法については、以前に記事にした。
その時にいくつか誤りがあったので訂正しておく。

畳み込み(cudnnConvolutionForward)などのAPIには、alphaとbetaというスケールパラメータがあるが、前回は__halfにしていたがfloatにする必要があった。
また、BatchNormalizationのパラメータは、入力と出力のディスクリプタがFP16であってもパラメータはすべてfloatにする必要があった。

以上を修正することで、正しく動作するようになった。

精度比較

FP32とFP16の方策ネットワークの出力を比較すると、

FP32 -23.4633 -18.3698 -26.9312 -16.0737 ...
FP16 -23.6094 -18.4531 -26.8438 -16.0781 ...

出力の値に0.15くらいの誤差がある。

10万局面を推論した際のfloodgateの棋譜との一致率は、

FP32 FP16(TitanV) FP16(2080Ti)
指し手一致率 0.43797 0.43765 0.438
勝敗一致率 0.708297 0.707907 0.707756
評価値誤差(MSE) 0.0199915 0.0199959 0.0199937

多少の誤差があるが、実用上問題ないレベルのように思う。

速度比較

FP32とFP16で推論の速度を比較した。
TitanVと2080Tiの両方の結果を記載する。

前回、TitanVでTensorCoreを有効にするにはアルゴリズムをCUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMMに固定する必要があったが、今回は逆にCUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMMに固定するとFP32よりも遅くなったため、自動選択されたアルゴリズムをそのまま使用している。

  • 10万局面推論の平均時間
  • バッチサイズ128

単位:ms

FP32 FP16
TensorCore
なし
TensorCore
あり
TensorCore
なし
TensorCore
あり
2080 Ti 13.955 11.966 5.384 4.529
Titan V 11.963 6.557 5.890 4.938
1080 Ti※ 16.218 - - -

※CUDA 9.2+cuDNN 7.2

考察

2080 TiのFP16の効果

入力から出力まですべてFP16にして、TensorCoreを有効にすると、2080 TiがTitan Vを上回っている。

2080 Tiで、FP32(TensorCore無効)からFP16(TensorCore有効)の速度向上は、2.4倍になっている。

FP16のTensorCore無効からTensorCore有効の速度向上は、1.19倍であり、TensorCoreによる速度向上より、FP16化による効果が大きい。

FP16化の効果

2080 TiとTitan Vともに、FP32からFP16にするとTensorCoreを無効でも速度が向上している。
2080Tiでは2.02倍、TitanVでは1.58倍高速になっている。
CUDAコアで計算する場合でも、FP16化の効果は大きいと言える。

また、推論速度は、2080 TiがTitan Vを上回っている。
2080 Tiは、FP16により最適化されていると言えそうだ。

TensorCoreが使われているか検証

2080TiでTensorCoreが使われているかプロファイラで確認した。

  • CUDNN_TENSOR_OP_MATH有効時

f:id:TadaoYamaoka:20181008184359p:plain

  • CUDNN_TENSOR_OP_MATH無効時

f:id:TadaoYamaoka:20181008184609p:plain

CUDNN_TENSOR_OP_MATH有効/無効で、明らかに使用されているアルゴリズムが異なっている。
CUDNN_TENSOR_OP_MATHを有効にすると、turing_h1688cudnn_256x128_ldg8_relu_exp_small_nhwc_tnがトップに来ている。
これが、TensorCoreを使用するアルゴリズムであることを示す資料がないので推測になるが、状況からTensorCoreが使われていると思われる。

まとめ

入力から出力までFP16にすると、TitanVよりも2080Tiの方が、将棋AIの推論において、推論速度が速いことが確認できた。

GeForce RTX 2080 TiでTensorCoreを使う

GeForce RTX 2080 Tiを使って、ニューラルネットワークの推論でTensorCoreが使えるかを試してみた。

TitanVでは、TensorCoreを使うことで、将棋AIの推論が2.57倍高速になった
2080 Tiにも、TensorCoreが544個搭載されており、TitanVの640個より少ないが、CUDAコアで計算するよりも速くなると思われる。

環境構築

まず、CUDAのバージョンを2080 Tiに対応したバージョンにする必要があるので、最新10.0にバージョンアップした。
それに合わせて、cuDNNも最新の7.3.1にした。

ベンチマークプログラム

TitanVで測定したときと同じベンチマークプログラムを使用した。
将棋AIのResNet 10ブロックのニューラルネットワークの推論速度を測定する。
DeepLearningShogi/gpubenchmark.cpp at master · TadaoYamaoka/DeepLearningShogi · GitHub

測定結果

TitanVの速度も参考のため測定した。
また、1080 Tiの前回の測定結果も記載する。

結果は以下の通り。

  • 10万局面推論の平均時間
  • バッチサイズ128

単位:ms

GPU TensorCore無効 TensorCore有効 比率(無効/有効)
2080 Ti 13.955 11.966 1.17
Titan V 11.963 12.057 0.99
1080 Ti※ 16.218 16.201 1.00

※CUDA 9.2+cuDNN 7.2

考察

2080 Ti

2080TiでTensorCoreを有効にすると、1.17倍高速になっている。
前回、CUDA 9.2+cuDNN 7.2で、TitanVで測定した結果では、2.57倍だったのに対して高速化の幅は小さくなっている。

Titan V

一方、Titan Vでは、TensorCoreを有効にすると、速度が0.99倍と、前回の結果と大きく異なる結果になった。
原因を調べるため、プロファイラで使用されているアルゴリズムを確認したところ、前回は、volta_s884cudnn_fp16~がトップにあったが、volta_sgemm_128x64_nnになっていた。
f:id:TadaoYamaoka:20181001174513p:plain

cudnnGetConvolutionForwardAlgorithm_v7で、自動で選択されるアルゴリズムを調べたところ、前回は、すべての畳み込み層で、CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMMが使用されていたが、今回は25層のうち21層でCUDNN_CONVOLUTION_FWD_ALGO_WINOGRAD_NONFUSEDが使用されていた。

そこで、前回と同じアルゴリズムにするため、明示的にCUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMMが使用されるようにして再測定した。

GPU TensorCore無効 TensorCore有効 比率(無効/有効)
Titan V 11.963 6.557 1.82

結果、TensorCoreが無効の場合より、1.82倍高速になった。
前回は、TensorCoreを使った場合は、6.888msだったので、ほぼ同じになっている。
プロファイラで使用されているアルゴリズムを確認したところ、前回と同じvolta_s884cudnn_fp16~がトップになった。
f:id:TadaoYamaoka:20181001180228p:plain

また、前回よりTensorCoreが無効の場合の速度が速くなっていた。
CUDA 10.0+cuDNN 7.3.1でTensorCoreが無効のアルゴリズムが改善されているようだ。

なお、2080 Tiでは、自動で選択されるアルゴリズムは、すべてCUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMMとなっていた。

まとめ

GeForce RTX 2080 Tiでも、TensorCoreを使うことで推論の速度が改善されることが確認できた。

CUDA、cuDNNのバージョンによっては、自動で選択されるアルゴリズムではTensorCoreが使われない場合があることがわかった。
その場合、明示的にアルゴリズムを指定する必要がある。

cuBLASでTensorCoreを有効にする

前回cuDNNでTensorCoreを有効にして推論の速度を比較したが、cuBLASでもTensorCoreを有効にして推論の速度を比較してみた。

cuBLASでTensorCoreを有効にする

cublasSetMathMode(cublasHandle, CUBLAS_TENSOR_OP_MATH)

で、TensorCore演算を有効に設定する。

行列演算で、BLASの関数を拡張したcublas~Ex関数を使用する。
浮動小数の精度別の関数があるが、要素ごとの精度を指定できるcublasGemmExを使う場合は、

cublasGemmEx(handle, CUBLAS_OP_T, CUBLAS_OP_N, n, m, k, &alpha, W, CUDA_R_32F, k, x, CUDA_R_32F, k, &beta, y, CUDA_R_32F, n, CUDA_R_32F, CUBLAS_GEMM_DEFAULT_TENSOR_OP)

のように最後の引数に、CUBLAS_GEMM_DEFAULT_TENSOR_OPを指定する。

浮動小数の精度にCUDA_R_32F(単精度)を指定した場合も、TensorCoreが有効な場合は、内部でFP16に変換されて計算される。

参考:
cuBLAS :: CUDA Toolkit Documentation

速度比較

TensorCore有効/無効で、推論の速度を比較した。
TensorCoreを搭載していないGeforce 1080 Tiでも測定した。

将棋AI用の10ブロックのResNetのモデルで、cuDNNのTensorCoreは有効にしている。
cuBLASは、価値ネットワークの出力の全結合層(中間ノード数256)で使用している。

結果は以下の通り。

10万局面推論の平均時間
バッチサイズ128
単位:ms

GPU 有効 無効 比率(無効/有効)
Titan V 7.041 6.850 1.03
1080 Ti 16.178 16.183 1.00
精度比較
無効 有効
指し手一致率 0.43799 0.43799
勝敗一致率 0.708287 0.708287
評価値誤差(MSE) 0.0199915 0.0199915

考察

cuBLASでTensorCoreを有効にしてもほとんど速くなっていない(誤差の範囲)。
1080 TiではTensorCoreを有効に設定しても、結果は変わっていない。

TensorCoreが有効になっているかプロファイルで確認したところ、
f:id:TadaoYamaoka:20180929122105p:plain
TensorCore無効/有効で同じアルゴリズム(volta_sgemm_32x32_sliced1x4_tn))が選択されていた。

CUBLAS_GEMM_DEFAULT_TENSOR_OPを指定しても、必ずTensorCoreが使われるわけではないことが確認できた。
FP16の場合は、TensorCoreが使われる可能性があるので、次はFP16にして試す予定。

追記

FP16にして測定してみた。
前回書いた方法でFP16化したが、計算結果がおかしくなってしまった。
バグとれていないので結果があやしいが、とりあえず速度だけ測定すると、8.438msとなり、FP32よりも遅くなった。
プロファイラで確認すると、
f:id:TadaoYamaoka:20180929133854p:plain
TensorCoreが使われているように見えない。
FP16化は苦労しそうなので、とりあえずFP32でもcuDNNでTensorCoreが使えているので調査はここまでにする。