例1:単精度のa*X plus Yの計算
例1の単精度のa*X plus Yの計算は要素ごとにベクトルの和を計算する。CUDAでは10.1ライブラリのcublasSaxpyでは32bit、あるいは64bitのグローバルメモリのload/store命令を使っている。
しかし、T4 GPUはスケジューラ当たり4個のLSUを持っている。なお、V100 GPUの場合はスケジューラ当たり8個のLSUと2倍のメモリバンド幅を持っている。そして、T4 GPUはSMあたり1024スレッドをサポートしている。これもVoltaは2048スレッドをサポートしている。
しかし、ブロック数やスレッド数を増やすだけでは、Turing GPUの持つメモリバンド幅を使い切ることは難しい。
これを128bitのベクタ化したメモリアクセス命令に替えると1命令でアクセスできるメモリ量を増やすことができる。
次の図の左側のコードではld.global.v4.f32という32bitのロード命令を使っているが、右側のコードではLDG.E.128.SYSで128bitのベクタのロード命令を使っている。このため、命令当たりのメモリバンド幅は4倍になっている。
次の図は、Saxpyの実行時間を示すグラフで、横軸はアレイのサイズ(KiB)である。この図からわかるように、青線のベクタロード命令を使うと、20KiBから2000KiBの範囲サイズのアレイではほぼ2倍の性能となっている。なお、2000KiB以上では差が減少しているが、その原因についての考察は書かれていない。多分、TLBミスとそれに伴うテーブルウオークが影響していると思われる。
例2:マトリクスの積和の計算
例2は、C+=ABというマトリクスの積和の計算である。それぞれのスレッドは8×512のA_sliceと512×8のB_sliceを掛けてC_tileの値を計算する。この行列積和の計算は多くのワークロードで最も計算負荷の高い処理となっている。
レジスタファイルのアクセスのボトルネックを考えるとき重要なことは、
- 命令は入力オペランドを必要とし、それらはレジスタファイルのポートを通してアクセスされる
- ポートが不足するとその命令の実行はストール(停止)する
- ポートアクセスを緩和するにはレジスタ再利用キャッシュを活用すべき
- コンパイラはレジスタ再利用キャッシュを使おうとするが、必ずしも最適な使い方になっていない
- NVIDIAのライブラリは、最適になっていない場合は人手で最適化している。つまり、我々も同じことをやれば良い。
NVIDIAのドキュメントには書かれていないのであるが、2018年に公開されたCitadelの論文では、NVIDIAの(少なくともVolta)GPUはレジスタ再利用キャッシュという構造を持っている。レジスタ再利用キャッシュは、メインのレジスタファイルに比べて小さなレジスタファイルで、メインのレジスタファイルとは独立にアクセスできるポートを持っている。
メインのレジスタファイルをアクセスすると、そのデータはレジスタ再利用キャッシュに記憶され、その値は、メインのレジスタファイルのポートが使えなくてもレジスタ再利用キャッシュから読み出すことができる。つまり、メインのレジスタファイルのポート不足を補うことができるという仕掛けになっている。
次の図の左側のコードで最初の命令は浮動小数点の積和演算で、入力はR12、R80、R16で積和の結果はR16に格納される。2番目の命令は同じ浮動小数点の積和演算で、入力はR80.reuse、R13、R17で積和の結果はR17に入れられる。この2番目の命令ではR80.reuseが指定されており、この場合はレジスタ再利用キャッシュから読まれ、メインのレジスタファイルの読み出しポートを使用しない。
メインのレジスタファイルの読み出しポート数が2つであるとすると、最初の命令は1サイクルでは全部の入力オペランドは読めず、オペランドの読み込みに少なくとも1サイクル余計に必要とする。一方、2番目の命令ではR80.reuseを入力に使っているので、メインのレジスタファイルから読むのはR13とR17の2つで、1サイクルで読むことができる。
右側のコードは、左側のコードをCitadelが最適化したもので、reuseと書かれた入力が多くなっている。結果として、Citadelのコードでは全部の命令が1サイクルで入力オペランドが読めるが、元の左側のコードでは4つの命令でオペランドの読み込みに追加のサイクルを必要とする。
このため、T4 GPUでの実行では、右のコードは左のコードと比べて12%高い性能が得られた。そして、この性能はNVIDIAが最適化したcuBLASの性能と一致したという。
GPUのメーカーは、このようなアーキテクチャを公表しないので、普通のGPUプログラムの開発者はこのような最適化を行うことはできない。GPUのアーキテクチャを理解するためには大変な努力を必要とするが、彼らがそれを行ったので、皆がその成果を利用できるようになる。この成果はarxive.orgに公開する準備を進めているという。
(次回は4月10日に掲載します)