GPUでスレッドを同期させる手法

スレッドの同期の必要性

GPUでは多数のスレッドを並列に動作させて処理を分担させる。このような場合、しばしば、すべてのスレッドの処理の終了を待ち合わせることが必要になる。例えば、各スレッドで何らかの処理を行い、その結果の合計を求めるという処理の場合は、一部のスレッドの処理が終わらないうちに、合計を求める処理を始めてしまうと正しい合計とならない。したがって、合計の処理を始める前に、すべてのスレッドの処理の終了を確認する必要がある。

NVIDIAのGPUでは、1つのワープに含まれる32スレッドは1つの命令列を実行するので、すべてのスレッドでの各命令の実行は同じタイミングで行われる。したがって、ワープの中では、どれか1つのスレッドでも終了すれば、他のスレッドも同じ命令まで実行されていることは保証される。

しかし、同一ブロック内でも、異なるワープの命令は切り替えながら順に実行されていくので、1つのワープでは処理が終了するところまで実行が進んでいるが、他のワープでは、(特に、キャッシュミスなどの時間がかかる処理が入ったりすると、)まだ、そこまで命令の実行が進んでおらず、処理が終わっていないということが起こり得る。なお、これはGPUだけの特殊事情ではなく、CPUでもマルチコア、マルチスレッドの場合は同じことが起こる。

このため、すべてのスレッドが特定の位置まで進んだことを確認する必要がある場合には、特別な操作が必要である。この確認を行うことを、スレッド間の同期をとると言う。

CUDAの「__syncthreads( )」

CUDAでは、「__syncthreads( )」という組み込み関数があり、この関数を呼び出すと、すべてのスレッドが同期するのを待ち合わせて、この関数からリターンする。__syncthreads( )の呼び出しは、先に終わったスレッドも、すべてのスレッドが終わるまで待ち合わせることになるので、性能的には、できるだけ挿入箇所を少なくする方が良い。一方、余計に入れても、性能的には低下するが、コード動作の正しさは変わらないので、安全サイドで入れるというのも有りである。

排他的なアトミックメモリアクセス

__syncthreads( )関数を作るには、例えば、スレッドの開始時点で、共通変数の値を0にしておき、終了したスレッドは、その共通変数の値を+1することにし、その値が同期すべきスレッド数になったら、すべてのスレッドが共通変数を+1するところまでは実行が進んだことが保証される。つまり、合計を求める処理は、共通変数の値をモニタし、それが同期すべきスレッド数になったら、__syncthreads( )関数からリターンして、次の処理、例えば合計を計算する処理を始めれば良い。

しかし、多数のスレッドがこの共通変数を読み、+1して書き戻すという処理を行う場合、スレッドiが共通変数を読み、+1して、その次の結果の書き戻しを行う前に、他のスレッドjが共通変数を読んでしまうということが起こり得る。そうすると、スレッドiとスレッドjは同じ値を読み、それを+1して書き戻すので、一方のスレッドの更新が上書きされて無視されてしまう。結果として全スレッドが処理を終わっても共通変数は同期すべきスレッド数にならず、プログラムがハングしてしまう。

この問題が起こらないようにするには、スレッドiが共通変数のメモリアクセスを開始したら、+1して書き戻しが終わるまで、他のスレッドが共通変数をアクセスできないようにすれば良い。このようなメモリアクセスを「アトミックアクセス」と言い、CPUと同様に、GPUでもアトミックなメモリアクセスができるようになっている。カウントを行う共通変数に対しては、アトミックなメモリアクセスを使えば、スレッドiの書き戻しが終わる前に、スレッドjが共通変数を読むということは起こらないので、問題は発生しない。

このような共通変数は、同一ブロック内のスレッドの同期の場合は、シェアードメモリ上に設ければ、高速でアクセスでき、効率が良い。一方、グリッドレベルの全スレッドの同期を必要とする場合は、グローバルメモリ上に共通変数を設ける必要がある。

なお、前世代のKepler GPUと比べると、Maxwell GPUではシェアードメモリのアトミックアクセスの性能が大幅に改善されており、頻繁に同期を必要とするプログラムでも同期のオーバヘッドが小さくなっている。