C言語をGPU用に拡張した言語「CUDA」

CUDAでのカーネルの実行指示

CUDAは、NVIDIAがC言語をGPU用に拡張した言語で、特徴的な構文が、

Mykernel<<<numBlocks、threadsPerBlock>>>(引数リスト);

というGPUで実行する関数(カーネルと呼ぶ)を呼び出す構文である。Mykernelは呼び出す関数の名前で最後に(引数リスト)が付くのは通常のC言語の関数呼び出しと同じであるが、CUDAでは多数のスレッドを起動する<<< >>>の部分が付加されている。この構文でnumBlocks個のスレッドブロックが起動される。これを「グリッド(Grid)」と呼ぶ。そして、各スレッドブロックにはthreadsPerBlock個のスレッドが含まれるので、結果として、numBlocks × threadsPerBlock本のスレッドが実行されることになる。なお、CUDAでは変数のタイプも拡張されており、これらの変数は3要素のuint3型とすることができる。次の図はnumBlocksが3、2、1、threadPerBlockが4、3、1の場合のグリッドとスレッドブロックを示している。

Kepler GPUではnumBlocksのX方向は最大 、YとZ方向は最大までのサイズを指定できるので、数字の上ではギガ以上のスレッドブロックを扱える。また、threadsPerBlockはX、Y方向は最大1024、Z方向は最大64までの数を指定できる。

実行するスレッドの数を指定するだけなら、numBlocks、threadsPerBlockと2つの指定に分ける必要はないのであるが、この2つはハードウェア的には大きな違いがあり、後述のように性能にも大きな影響を及ぼすパラメタとなっている。

このようにカーネルの実行を指示された場合、ギガスレッドエンジンは、グリッドに含まれるスレッドブロックをSMに順に割り付けていく。後述のようにSMの資源が許す範囲で、複数のスレッドブロックを1つのSMに割り付ける。しかし、グリッドのサイズが大きい場合は、すべてのスレッドブロックをSM群に割り付けることはできないので、どこまで割り付けたかを記憶しておき、割り付けたスレッドブロックの実行が終わりSMの資源が空くと、残りのスレッドブロックを割り付けるというようにして処理を進めていく。

1つのスレッドブロックの中の全てのスレッドは1つのSMで実行され、2つ以上のSMに分割されることはない。つまり、1つのスレッドブロックに含まれるスレッドは、割り付けられたSMの内部のシェアードメモリを共用する。従って、同一スレッドブロック内のスレッド間では、シェアードメモリを経由してデータのやり取りが可能で、これは比較的高速に実行できる。これに対して、別のSMで実行されるスレッドとの間のデータ交換はDRAMのグローバルメモリ経由となるので400～800サイクルのメモリアクセスが2回必要となり、時間が掛るということになる。

CUDAプログラムでは、CPUメモリからGPUメモリへのデータのDMA転送、いくつかのカーネルの実行、そして結果のGPUメモリからCPUメモリへのDMA転送という流れが一般的である。この一連の処理は、前の処理が終わると次の処理が始まるというインオーダ実行であり、「ストリーム」と呼ばれる。CUDAでは相互に非同期で並列に実行される複数のストリームを定義することができる。この場合、1つのストリーム内では実行はインオーダであるが、別のストリームを並列に実行することができ、別のカーネルのグリッドを同時にSMに実行させることが可能になる。1つのグリッドだけではSMに十分な数のワープを確保できない場合は、このように複数のグリッドを実行させることにより、SMが担当するワープ数を増やして、GPUハードウェアの利用率を改善することができる。