TadaoYamaokaの日記

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

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)で調べてみた。
f:id:TadaoYamaoka:20180923163529p:plain

volta_s884cudnn_fp16_256x64_ldg8_relu_exp_small_nhwc2nchw_tn_v1
という項目があり、以下のページによるとs884cudnnはTensorCoreのオペレーションを表しているようだ。
Mixed precision training using float16 — mxnet documentation
実際にTensorCoreが使われているようである。