GPUの実効性能を上げるDynamic ParallelismとHyper-Q

NVIDIAのKepler GPUでは、Dynamic ParallelismとHyper-Qという新機能が搭載された。従来、この機能が使えるGPUは、Top500 1位のTitanスパコンなどに使われているK20とK20Xという科学技術計算用のアクセラレータボードだけであったが、2013年2月に発表されたコンシューマ向けのGeForce GTX Titanでもサポートされ、1000ドル以下のGPUボードでDynamic ParallelismとHyper-Qが使えるようになった。

これは比較的安価なGTX Titanを搭載するPCやワークステーションで、Titanスパコンなどのソフトウェアの開発が可能になり、また、大規模なスパコンで開発された構造解析や薬剤シミュレーションなどのプログラムが(解ける問題規模は小さいが)ワークステーションで実行できるようになることを意味している。

GPUでの計算プログラム

NVIDIAのGPUで計算を行う場合、CUDAという言語が使われるのが一般的である。2次元の行列の和を求めるCUDAプログラムは次のようになる。なお、この例は、NVIDIAのCUDA C Programming Guideからとったものである。

// Kernel definition
 __global__ void MatAdd(float A[N][N], float B[N][N],float C[N][N])
{ 
int i = threadIdx.x;
int j = threadIdx.y; 
C[i][j] = A[i][j] + B[i][j];
}

これはMatAddという計算カーネルで、GPU上で実行されるプログラムである。このカーネルは A、B、Cという3つのN行、N列の行列を引数としている。threadIdx.xというのはシステム定義の変数で、この2次元に配列されたスレッドのx方向の番号、threadIdx.yはy方向の番号である。これをiとjという変数に代入し、C[i][j] = A[i][j] + B[i][j]で行列AとBのi、j要素の和をCのi、j要素に格納している。

汎用のマイクロプロセサの場合は、スレッドの生成や切り替えのオーバヘッドが大きいので、各スレッドはある程度まとまった仕事をさせるのが普通であるが、CUDAでは、この例のように1つのスレッドは行列の1つの要素の値を計算するというのが一般的なプログラムの書き方である。

そして、CPU上で実行され、この計算カーネルを呼び出すメインプログラムは次のようになる。

int main() 
{
    ...
    // Kernel invocation with one block of N * N * 1 threads
    int numBlocks = 1;
    dim3 threadsPerBlock(N, N);
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
}

MatAdd<<< >>>の部分が上に書いたカーネルを呼び出す部分で、(A,B,C)はカーネルに渡す引数である。CUDAで特徴的なのが<<<　>>>で、ここで指定した数のスレッドをGPUに並列に実行させる。numBlocksはブロックの数で、この例では2行前で1が代入されている。そしてthreadsPerBlockは1つのブロックに含まれるスレッドの数で、dim3と定義されているので3次元(3要素)の変数であるが、ここでは2次元分だけを使い、その値はN,Nとなっている。

つまり、MatAdd<<<　>>>の行は、x方向N個、y方向N個の行列の各要素に対応し、全体ではN×N個のスレッドを含むブロックを1個起動するという意味である。

このカーネルとブロック、スレッドの関係を示すのが次の図である。

このようにCUDAでは最大3次元のスレッド配列であるブロックを2次元に並べることができるので、各次元の要素数にもよるが、膨大な数のスレッドを1つの文で起動することができる。なお、CUDAではブロックを並べたものをグリッド(Grid)と呼び、MatAdd<<<　>>>の行はグリッドの実行をGPUに依頼する文となっている。

NVIDIA GPUでのスレッド処理

NVIDIAのKeplerアーキテクチャのGK110チップの演算を行うSMXと呼ぶブロックは次の図のようなコンポーネントを含んでいる。そして、GK110チップには、このSMXが15個と、CPUと接続するPCIexpress3.0インタフェース、L2キャッシュ、6チャネルのGDDR5メモリインタフェースが集積されている。なお、Titanスパコンなどの使用されているK20Xでは15個全部のSMXが使われているが、K20とGeForce GTX Titanでは1個はスペアで残りの14個のSMXを使用している。

そして、1つのブロックに含まれる全スレッドは必ず1つのSMXで実行され、2つ以上のSMXに分割して実行されることはない。一方、1つのSMXで実行されるブロックは1つとは限らず、複数の異なるブロックが同じSMXで実行されうる。

NVIDIAのGPUではCUDAを使って科学技術計算を行う以前から、32スレッドを単位として並列に実行する方式が採られており、この32スレッドのまとまりをワープ(Warp)と呼んでいる。なお、このWarpはハイパードライブで宇宙の離れた2点間を瞬時に移動するという意味ではなく、Threadの横糸と直交する縦糸という意味である。

1つのWarpに含まれる32のスレッドは同一のカーネルのスレッドであり、すべて同じ命令を実行するが、使用するレジスタファイルはそのスレッドに割り当てられた占有部分となり別々のデータを処理する。また、thredIdxの値もスレッドの位置に対応する値が供給される。

1つのブロックに含まれるスレッドはWarpという単位に分割され、Warpスケジューラで命令発行される。そして、この図のように、Warpスケジューラは時間単位ごとに異なるWarpの異なる命令を発行していく。Warpスケジューラは、毎サイクル、Warpの2つ命令を発行することができ、1つのSMXには4つのWarpスケジューラがあるので、毎サイクル8命令を発行することができるようになっている。

このようにWarpを切り替えて命令発行すると、例えば最初に実行を開始したWarp 8の命令の演算やキャッシュアクセスにある程度の時間が掛っても、次のWarp 8の命令が出てくる時には処理は終わっており、これらの命令の実行レーテンシを隠すことができる。

また、1つのSMXでは同時に発行できるのは最大4Warpであるが、このように時間ごとにWarpを切り替えることにより、最大64Warp(2048スレッド)を(時分割で)並列に処理することができる。なお、SMXのレジスタファイルは64Kエントリと非常に大きいが、2048スレッドで共有すると1スレッドあたりでは32エントリであり、汎用プロセサのレジスタ数と大きな違いは無い。

WarpとSMXのCUDAコアの対応がどうなっているかは公表された情報はないが、ブロック図で縦方向に並んだ16個の実行ユニットが最初サイクルでWarpの前半を実行し、次のサイクルでも同じ命令を実行してワープの後半を実行しているのではないかと思われる。