Voltaの演算性能

次のグラフは、CUDA9.0のcuBLASを使って行列の乗算を行った場合の性能をプロットしたものである。グラフの横軸は行列のサイズで、縦軸はTFlops値である。3本の折れ線は、下から順に、倍精度浮動小数点演算、単精度浮動小数点演算、Tensorコアを使う混合精度浮動小数点演算の場合である。

行列のサイズがある程度以上大きい場合は、倍精度では7TFlops程度、単精度では14TFlops程度の性能が得られており、これは理論ピーク性能とほぼ一致している。混合精度の場合は、ギザギザした線になっているが、70TFlops強の性能が得られている。理論ピーク性能は113TFlopsであり、cuBLASの使用でピークの70%程度の性能が得られている。

  • Voltaの浮動小数点演算の性能

    Voltaの浮動小数点演算の性能。行列のサイズが2Kを超えるとほぼ一定の性能になり、倍精度では7TFlops、単精度では14TFlops、混合精度を使うTensorコアの積和演算では70TFlops強の性能が得られている

データキャッシュのレーテンシとバンド幅

次の図の2つの棒グラフは、左がL1データキャッシュのヒットの場合のレーテンシ、右がL1データキャッシュのバンド幅を示している。VoltaはこれまでのNVIDIAのGPUの中でも一番短い28サイクルのL1データキャッシュレーテンシと一番高い108Byte/cycleというバンド幅を実現している。

VoltaのL1データキャッシュはLRUではなく、各セットから直前にロードされたキャッシュラインを入れ替えの候補としている。頻繁に入れ替えが起こるデータをこの部分でカバーし、入れ替えが頻繁でない行列のようなデータがキャッシュから追い出されるケースを減らしている。

余談であるが、富士通のスパコン用のCPUは、キャッシュを頻繁な入れ替え用のWayと入れ替えの少ないWayに分割し、それぞれの変数をどちらに入れるかをコンパイラに指示できるようになっている。構造やサイズは違うが、入れ替えに関してはVoltaも同じ考えを取り入れている。

面白いのは、Voltaに次いでレーテンシやバンド幅性能が高いのはKepler K80で、MaxwellやPascalのL1データキャッシュは性能が低いという点である。MaxwellとPascalはL1データキャッシュとシェアードメモリを分離した構造であるが、KeplerはVoltaと同じく両者を一体とした構造になっており、一体型の構造を使う方が性能が高いという結果になっている。

  • ></span>L1データキャッシュのヒットレーテンシとバンド幅

    L1データキャッシュのヒットレーテンシ(左)とバンド幅(右)。一体型の構造になっているV100と3世代前のKeplerの性能が高い

Voltaは命令レーテンシを2/3に短縮

次の図は、PascalとVoltaの代表的な命令の実行レーテンシを示している。Pascalでは大部分の単純な演算命令のレーテンシは6サイクルであったが、Voltaでは、これが4サイクルに短縮されている。このため、P100とV100 GPUのクロックはほぼ同じであるが、命令のレーテンシがネックになる処理では、Voltaの方が1.5倍性能が高くなっていると言える。

なお、PascalのIMULとIMADが～86サイクルと非常に長い時間が掛かっているが、これはPascalはこれらの命令を直接実行できるハードウェアを持っておらず、エミュレーションで実行しているからである。

  • PascalとVoltaの代表的な命令の実行レーテンシ

    PascalとVoltaの代表的な命令の実行レーテンシ。大部分の演算命令はPascalでは6サイクル掛かっていたが、Voltaでは4サイクルに短縮されている

Tensorコアの計算の仕組み

Volta V100 GPUの目玉とも言えるTensorコアは16行16列の行列同士の積を計算する。

次の図のように、16行16列の行列を8個のグループに分けて考え、グループ0の各行はスレッド0～3の1つのスレッドが分担し、グループ2の各行はスレッド4～7が分担するというやり方で1つのワープの32スレッドで計算を行う。

ただし、この行列積を計算するwmma::mma_syncというハードウェア命令を持っているわけではなく、NVCCコンパイラがwmma::mma_sync を16個のHMMA命令(FP16の積和演算命令)に変換している。

ここでは、青いボックスで囲まれた4命令をセットと呼び、各セットは次の2枚目の図のように行列aと行列bの異なる部分を計算するが、結果は行列accの同じところに足し込んでいく。

それぞれのセットの中では異なるSTEPフラグを持つ命令は行列accの異なる場所を計算している。なお、オレンジの16要素は1巡目に計算される部分で、青の16要素は2巡目に計算される部分を示している。

「Voltaで新設されたTensorコア」の記事で筆者が推定した計算方法は、NVIDIAの実装とは異なっていた。全体の計算量は同じであり、複雑に見えるこの実装をNVIDIAが使ったのは、必要なデータのロードやReuse cacheがどれだけ有効に使えるかなどの点でこの実装の方が有利であるのかもしれない。

  • Tensorコアの処理の説明図
  • Tensorコアの処理の説明図

  • Tensorコアの処理の説明図

(次回は5月31日に掲載します)