将棋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が使われているようである。
