TadaoYamaokaの日記

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

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の推論において、推論速度が速いことが確認できた。