CUDA の Tensor Core を使ってみた

折角 TITAN V を購入したので、SM 7.0 以上で利用できる
Tensor Core を使った行列計算テストプログラムを
書いてみました。

CUDA Toolkit のサンプルとして Tensor Core を
使ったプログラムがありますが、余計な最適化処理を
していてちょっと分かりづらい。

さらに2のn乗(n >= 4)でないと動作しないコードだった
ので、任意の行列に対応するようにしました。

で、その効果ですが、やや微妙です。。。

start wmma version...
calc_fm elapsed time without TensorCore:7.1383ms
calc_fm elapsed time without TensorCore:4.08064ms
calc_fm elapsed time with TensorCore:3.6567ms
start CPU gold...
calc_fm_gold elapsed time:43.6952m

MNIST を VGG とかで解析する際の第1段めを想定した実験
ですが、効果は10%程度となりました。っていうか、Tensor Core
を使わなくても TITAN V の CUDA Core が速すぎるというか。

WMMA を使うにあたっていくつか TIPS がありました。

1. 行列の内積を求める際は row_major を使う。
   サンプルでは matrix_a に対して col_major を指定しています
   が、row_major にしないと正しく積が計算されません。
2. __syncthreads() を呼ばなくてもいい。
   wmma ロジック内部でスレッド同期してから返される
   ようなので、__syncthreads() 等のスレッド同期を外部で
   行う必要はありませんでした。
3. タイル数をむやみに増やしても効果がない。
   nVidia のサンプルでは全ての部分行列（タイル）を読み込んで
   高速化を図っていますが、行列をキャッシュして使い回す以上の
   意味はなく、それ以外の積和演算のパフォーマンスへの
   インパクトがありません。
   これは wmma のアクセス関数が *_sync と同期タイプしか
   提供されていないためだと思われます。
   （正直、どうやったら 640 コア使えるのか分かりません。）
4. 部分行列の総和は matrix_c を読み書きする。
   accumulator はload/storeが可能です。
   というか、この機能がないとwmmaの外で総和を計算
   することになり、若干不便です。（私が複数タイル用に
   書いたようなコードを埋め込む必要が出てきます。）
5. 行列の要素数によっては効果がないことがある。
   一般に、要素数が多いほど効果は上がるのですが、
   n, m, k の関係によっては Tensor Core を使った方が
   遅くなる現象が見られました。発生条件は調査中です。