Kepler GPUにおける命令の発行方法

Kepler GPUの命令の発行

Kepler GPUでは、ギガスレッドエンジンからSMに割り付けられたスレッドブロックのワープ群を4つのワープスケジューラに割り振る。各ワープスケジューラは最大16個のワープを分担し、これらのワープを並列に処理することができるようになっている。SMは命令キャッシュを持っており、ワープスケジューラは命令キャッシュ経由でDRAMに格納されたカーネルプログラムのGPU命令を読み出して命令の解釈を行う。

各スケジューラは担当している16ワープそれぞれの命令ポインタが指す次に実行する命令に対して、CUDAコアやDP Unit、LD/ST、SFUなどの実行ユニットが使用可能であり、かつ、 32スレッド全部の入力データが揃っていることをチェックし、これらの条件が揃い実行可能な命令を選んで32レーンの実行ユニットにその命令を発行する。つまり、ワープの実行ユニットの中の1つのレーンでも入力待ちになっているスレッドがあると、ワープ全体の命令が発行できないことになる。また、実行可能な命令が無い場合は、どのワープの命令も発行できず、実行ユニットが遊んでしまうということが起こる。

それぞれのワープスケジューラは2つのディスパッチユニットを持っており、実行可能な命令の次の命令も実行可能である場合は、2つの連続する命令を同時に発行することができる。

この図のように、ワープスケジューラは命令を発行するワープをサイクルごとに切り替え、分担するすべてのワープの命令が公平に実行されるようにしている。また、ワープを切り替えて実行することにより、この図に見られるようにワープ8の命令11、12の発行から、次にワープ8の命令13、14の発行までには他のワープの命令を実行するサイクルが入り、命令13、14が命令11、12の演算結果を使う場合の演算レーテンシを隠すことができる。

なお、この図では常に右側のディスパッチユニットも命令を発行しているが、命令12が命令11の結果を使う場合や、実行ユニットが不足する場合など、後続の命令が発行できないケースがある。その場合は、最初の命令だけが発行され右側の命令スロットは空くことになる。また、DP UnitはCUDAコアの2倍の64bitのデータを扱い2命令分のレジスタファイルのポートを使うので、他の命令と同時に発行することはできないと思われる。

SMには4つのワープスケジューラがあり、それぞれ最大2命令を発行することができるので、1サイクルに最大8命令を発行することができるが、単精度浮動小数点や整数演算を行うCUDAコアは6命令分、DP Unitは2命令分、LD/STとSFUは1命令分しかハードウェアがなく、同時に発行する命令が使う実行ユニットがこの範囲内となっていないと命令は発行できない。

また、演算には10～20サイクル、メモリアクセスには400～800サイクルかかるので、その結果を入力として使う命令は、演算やメモリアクセスが終わらないと入力待ちで命令を発行できないということが起こる。

ワープ内では命令は順に(インオーダ)実行されて行くが、このようにダイナミックにワープの実行順序が変わるので、ワープ間では命令の実行順序は保証されない。このため、他のワープの処理結果を使う場合は__syncthreads( )を実行し、データ生成側のワープの終了を待って、処理結果を利用する必要がある。

スレッドブロックのサイズはどう選ぶのか

ワープスケジューラは、1つのワープの命令が発行できなくなっても、そのスケジューラが分担しているワープ群の中から実行可能な命令を持つワープを探して命令を発行するので、直ちに命令の発行が止まるわけではない。しかし、SMに割り当てられているワープ数が少ないと、実行可能な命令を持つワープが見つからないということが起こる可能性が高くなる。

KeplerのSMは4個のワープスケジューラを持っており、命令の発行を止めないようにするには、少なくとも実行条件の揃った4つのワープが存在する必要がある。そして、演算の結果が得られるまでのレーテンシを10サイクルとすると、4×10=40ワープ程度が必要となる。ただし、直前の命令の演算結果を使わない命令が相当数ある場合は、これより少ない数のワープでも命令の発行は止まらず、NVIDIAは、SMあたり32ワープ程度を目安としてプログラムを作り、最適化の過程でワープ数を増減して性能が高くなるように調整することを勧めている。

一方、LD/ST命令がL2キャッシュをミスしてDRAMにアクセスする場合は、完了までに400～800サイクルという長い時間が掛るので、そのワープの命令発行は、メモリアクセスが終わるまでお休みとなり、残ったワープの命令の発行を行うことになる。

スレッドブロックの中のスレッドは、dim(X)＊｛dim(Y)＊Z＋Y｝＋Xで計算される一連の番号が付けられ、最初の32個が最初のワープ、次の32個が2番目のワープというようにワープに分割される。なお、dim(X)、dim(Y)はスレッドブロックのX、Y方向の要素数、X、Y、Zはそのスレッドのスレッドブロック内の座標である。

ワープは、必ず32スレッド単位で実行されるので、スレッドブロック内の総スレッド数が32で割り切れない場合は、最後のワープにダミーのスレッドを加えて32スレッドにして実行する。このダミーのスレッドは有効な計算には貢献しないので、スレッドブロックの総スレッド数は32の倍数にすることが望ましい。特に、総スレッド数が少ない場合には、この影響が大きく、次の図に示す80スレッドの場合は、最初の2つのワープはすべて有効なスレッドであるが、3番目のワープは16しか有効なスレッドが入っておらず、 GPUのピーク性能の80/96 ≈ 83.3%が達成可能な性能の上限となってしまう。

つまり、スレッドブロックに含まれるスレッドの数は32の倍数が望ましく、SMが担当するワープ数を大きくするためには多くのスレッドを含むことが望ましい。しかし、スレッド数をあまり大きくすると、レジスタやシェアードメモリが不足したり、SMへのスレッドブロックの詰め込みが難しくなったりするので、これらの制約も考慮してスレッドブロックのサイズを決めることになる。