GeForce RTX 2080 TiでTensorCoreを使う(その２ 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	6.895	5.800
Titan V	11.963	6.557	7.543	6.323
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有効時

• CUDNN_TENSOR_OP_MATH無効時

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になっていた。

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~がトップになった。

また、前回より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が有効になっているかプロファイルで確認したところ、

TensorCore無効／有効で同じアルゴリズム(volta_sgemm_32x32_sliced1x4_tn))が選択されていた。

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

追記

FP16にして測定してみた。
前回書いた方法でFP16化したが、計算結果がおかしくなってしまった。
バグとれていないので結果があやしいが、とりあえず速度だけ測定すると、8.438msとなり、FP32よりも遅くなった。
プロファイラで確認すると、

TensorCoreが使われているように見えない。
FP16化は苦労しそうなので、とりあえずFP32でもcuDNNでTensorCoreが使えているので調査はここまでにする。

cuDNN/cuBLASのFP16対応

TitanVを使って、FP32からFP16にするとどれくらい推論が速くなるか試してみた。
また、Geforce 1080 Tiなどのゲーム用のGPUにはFP16のアクセラレータが搭載されていないが、FP16の演算は可能なため、FP32と速度が変わらなければメモリ転送の効率が上がる分だけ速くなると思われる。そのため、Geforce 1080 Tiでも測定してみた。

以前にMNISTデータセットを使ったResNetの推論をCUDA/cuDNN/cuBLASを使って実装した際のコードをFP16に対応させて比較した。

GitHub - TadaoYamaoka/mnistCUDNN at fp16

FP16対応

データ形式変換

FP32からFP16へのデータ形式の変換は、CUDAの__float2half関数で行える。
CUDA Math API :: CUDA Toolkit Documentation

演算のFP16化

cuDNNのAPIにはデータ形式を指定する引数があるので、CUDNN_DATA_HALFを指定する。
cuBLASでは、CUDA_R_16Fを指定する。

測定結果

MNISTデータセットで、1ブロックのResNetの推論の時間を、FP32とFP16で比較した。

結果は以下の通り。

20エポックの推論の合計時間
バッチサイズ：100
単位：秒

精度	GPU	時間	FP16/FP32
FP32	Titan V	1.257	0.997
FP16	Titan V	1.253
FP32	1080 Ti	1.484	32.223
FP16	1080 Ti	47.847

考察

TitanVではFP16にすると、処理時間99.7%と、ほとんど変わらなかった。
チャンネル数が1のMNISTデータセットでは、演算の処理時間はほとんど変わらなかったと思われる。

GeForce 1080 Tiでは、FP16にすると処理時間が32.2倍かかるようになり、極端に遅くなった。
cuDNNでの畳み込み演算で特に時間がかかっていた。
FP16アクセラレータがないGPUでもFP32とFP16の演算時間は変わらないと予測していたが、遅くなるという結果だった。

TitanVでFP16化の効果が測定できなかったので、次はネットワークの規模の大きい将棋AIの推論で比較してみる予定。

cuDNNでTensorCoreを有効にする

将棋AIの強化学習にTitan Vを使用しているが、今までTitan Vに搭載されているTensorCoreを使えていなかった。
cuDNN 7.1以前では、TensorCoreを有効にするにはプログラムをFP16に対応させる必要があった。
cuDNN 7.2で、FP32でもTensorCoreが使えるようになったので、どれくらい速くなるか試してみた。

TensorCoreを有効にする

畳み込み演算(Convolution)で、TensorCoreを有効にするには、ディスクリプタに対して、

cudnnSetConvolutionMathType(convDesc, CUDNN_TENSOR_OP_MATH_ALLOW_CONVERSION)

をコールすればよい。
GPUにTensorCoreがない場合も、TensorCoreが使われないだけなのでコールしても問題ない。

アルゴリズムの選択には、cuDNN v7のAPIを使う必要がある。

cudnnGetConvolutionForwardAlgorithm_v7(handle, xDesc, wDesc, convDesc, yDesc, 1, &returnedAlgoCount, &algo_perf)

自動的にTensorCoreが有効になるアルゴリズムが選択される。

ベンチマークプログラム

将棋AI用のResNet 10ブロックの推論の速度を以下のベンチマークプログラムで比較した。
DeepLearningShogi/gpubenchmark.cpp at master · TadaoYamaoka/DeepLearningShogi · GitHub

速度比較

CUDNN_TENSOR_OP_MATH_ALLOW_CONVERSIONあり／なしで、推論の速度を比較した。
TensorCoreを搭載していないGeforce 1080 Tiでも測定した。

結果は以下の通り。

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

単位：ms

	なし	あり	比率(なし/あり)
Titan V	17.697	6.888	2.57
1080 Ti	16.218	16.201	1.00

精度比較

TensorCoreを有効にすると、内部的にFP16に変換が行われる。
そのため、計算の精度が低くなる。
指し手と勝率の予測の結果にどれだけ影響があるか確認した。
floodgateの棋譜との指し手の一致率、勝敗の一致率、評価値の誤差で確認した。

	なし	あり
指し手一致率	0.43797	0.43799
勝敗一致率	0.708297	0.708287
評価値誤差(MSE)	0.0199915	0.0199915

考察

TensorCoreを有効にすると、2.57倍高速になった。
TensorCoreを搭載していない1080 Tiでは、CUDNN_TENSOR_OP_MATH_ALLOW_CONVERSIONを設定しても処理時間が変わっていない。
TensorCoreを有効にすることで、CUDAコアで計算するよりも高速になることが確認できた。

TensorCoreを使わない場合は、1080 Tiの方がわずかに速い。
ベースクロックが1080 Tiの方が高いためと思われる。
TensorCoreを使わなければ、Titan Vの性能は活かせないと言える。

計算の精度について、畳み込み演算が内部的にFP16に変換されるが、最終結果にほとんど影響していない。

cuBlasでもTensorCoreを有効にできるので、全結合の速度が改善できるはずなので次に試す予定。

今月発売されるGeforce 2080 Tiでも、TensorCoreが544個搭載されるので、機械学習目的でも使えそうです。
Titan Vの640個よりは少ないですが、コストパフォーマンスはよさそうですね。

2018/9/23 追記

実際にTensorCoreが利用されているかプロファイラ(nvprof)で調べてみた。

volta_s884cudnn_fp16_256x64_ldg8_relu_exp_small_nhwc2nchw_tn_v1
という項目があり、以下のページによるとs884cudnnはTensorCoreのオペレーションを表しているようだ。
Apache MXNet | A flexible and efficient library for deep learning.
実際にTensorCoreが使われているようである。

将棋AIの進捗 その25(自己対局による強化学習の経過)

前回からだいぶ期間が空きましたが、自己対局による強化学習で、教師ありで収束するまで学習したモデルより有意に強くすることができました。前回は、19イテレーションでほぼ互角の強さでしたが、38イテレーションまで自己対局を行うことで有意に強くなりました。

自己対局における探索延長

強化学習の方法は前回までの方法とほとんど同じですが、途中から探索延長を行うように変更しました。
探索延長は、固定回数シミュレーションの結果、1番目に訪問回数が多い手が2番目に訪問回数が多い手の1.2倍未満の訪問回数のときに、シミュレーション回数を1.5倍にするという実装です。
シミュレーション回数を増やすことで、生成される教師データの精度をあげることができますが、計算時間が増えてしまいます。
結果がわかりきった局面に対してもシミュレーション回数を増やすと無駄があるので、探索延長を行うことでより計算効率を上げることができます。

強さの測定方法

前回までは、平手の初期局面から対局を行っていましたが、多少ランダムは入れていますがソフト同士の対局ではゲームの進行が偏る傾向があるので、24手までは「やねうら互角局面集」を使うようにしました。

教師ありで収束するまで学習したモデルとの対局結果

100回対局の結果は以下の通りです。

対局数100 先手勝ち47(47%) 後手勝ち49(49%) 引き分け4
selfplay038
勝ち63(63%) 先手勝ち31(31%) 後手勝ち32(32%)
model110
勝ち33(33%) 先手勝ち16(16%) 後手勝ち17(17%)

selfplay038が自己対局で38イテレーション学習したモデルで、model110が、elmoで生成した11億局面を学習したモデルです。
100回対局で有意水準を5%とすると、臨界値は58.2%なので、有意に強くなっています。

floodgateの棋譜に対する一致率

イテレーション回数	policyの一致率	valueの一致率
19(前回)	0.4090266	0.68762606
38	0.437221587	0.706173778

イテレーションごとのfloodgate棋譜との一致率グラフ

GPSFishとの対局結果

GeForce 1080を搭載したノートPCでの対局結果です。
GPSFishとはほぼ互角になりました。教師ありで学習したモデルよりも勝率が上がっています。
eloレーティングで言うと、GPSFishのと対局を基準として前回から35.6の向上です。
GeForce 1080 TiだとGPSFishに勝ち越せそうですが、未測定です。

38イテレーション自己対局で学習したモデルとGPSFishの対局結果

対局数100 先手勝ち50(50%) 後手勝ち47(47%) 引き分け3
selfplay038
勝ち45(45%) 先手勝ち24(24%) 後手勝ち21(21%)
GPSfish 0.2.1+r2837 gcc 4.8.1 osl wordsize 32 gcc 4.8.1 64bit
勝ち52(52%) 先手勝ち26(26%) 後手勝ち26(26%)

教師ありで学習したモデルとGPSFishの対局結果

対局数100 先手勝ち50(50%) 後手勝ち47(47%) 引き分け3
model110
勝ち40(40%) 先手勝ち21(21%) 後手勝ち19(19%)
GPSfish 0.2.1+r2837 gcc 4.8.1 osl wordsize 32 gcc 4.8.1 64bit
勝ち57(56%) 先手勝ち29(28%) 後手勝ち28(28%)

自己対局にかかった時間と電気代

前回から、2GPUのPCを2ヵ月以上動かし続けました。
去年の同じ時期よりも、1月あたり200kWhくらい消費電力が増えて、電気代は8000円くらい増えました(´;ω;｀)

今後の見通し

20イテレーションで、レーティング35.6の向上なので、現在のトップが4300くらいなので、GPSFishと同じ3100とすると、あと1200必要です。
楽観的に線形にレーティングが伸びるとすると、674イテレーションは必要です（実際はlogの形なのでもっとかかる）。
1ヵ月10イテレーションとすると、67.4ヵ月＝5.6年以上かかる見込みです。

今後も継続するつもりですが、より短時間で学習できる方法を見つけたいところです。
もしくは他力本願でGPUの性能が数十倍になってくれることを祈っています。

Protocol BuffersをTensorBoardでグラフ表示

バリューネットワークにはプーリング層が有効らしく、AQでもバリューネットワークはプーリング層を使っているようなので、AQのニューラルネットワークの構成を調べてみた。

GitHubで公開されているソースでは、ニューラルネットワーク構成は、Protocol Buffersの形式で保存されていたので、TensorBoardで可視化してみた。

TensorBoardでProtocol Buffersのグラフを表示する方法

Protocol Buffers形式は、TensorBoardで直接開くことができない。
一旦ログ形式にする必要がある。

やり方はここに書かれていた方法を参考にした。
tensorboard: view graph from saved_model.pb file [feature request] · Issue #8854 · tensorflow/tensorflow · GitHub

import tensorflow as tf
from tensorflow.python.platform import gfile
with tf.Session() as sess:
    model_filename ='AQ/pb/vl.pb'
    with gfile.FastGFile(model_filename, 'rb') as f:
        graph_def = tf.GraphDef()
        graph_def.ParseFromString(f.read())
        g_in = tf.import_graph_def(graph_def)
LOGDIR='log'
train_writer = tf.summary.FileWriter(LOGDIR)
train_writer.add_graph(sess.graph)

logディレクトリにログ形式で保存されるので、

tensorboard --logdir log

のように実行する。

ブラウザでコントロールに表示されたURLを開くとニューラルネットワーク構成をグラフで表示できる。