シェアードメモリ経由のスレッド間通信

また、図3-19に見られるように、各々のスレッドは自分専用のレジスタファイルだけをアクセスして処理を行う構造になっており、他のスレッドのレジスタファイルをアクセスすることはできないので、スレッド間でのデータのやり取りはできない。

このため、スレッド間でデータのやり取りを行う場合には、シェアードメモリを使う。同一グループの32スレッドは同じシェアードメモリを使うので、スレッド1がアドレスAに書いたデータを、スレッド2がアドレスAを読むことで受け取ることができる。また、同一SM(Streaming Multiprocessor)で走るすべてのスレッドは同じシェアードメモリを使うので、これらのスレッドの間でのデータのやり取りもできる。

デバイスメモリを使えば、別SMで走るスレッドとの間でデータの受け渡しをすることができるが、デバイスメモリのアクセスはシェアードメモリのアクセスより時間が掛かるので、同一SMのスレッド間ではシェアードメモリを使う方が良い。

なお、同一グループの32スレッドは同期して実行されるが、デバイスメモリのアクセスの順序がプログラム順になることは保証されていない。また、別グループのスレッドの間では、命令の実行順序も保証されない。このため、スレッド間の同期が必要な場合は、_syncthreads( )関数の呼び出しを入れておくことが必要である。_syncthreads( )がある場合は、それ以前の命令によるすべてのメモリアクセスが完了してから、それ以後の命令のメモリアクセスを開始するので、メモリアクセスの完了の順序が前後してしまうことはない。ただし、このような同期を行うと性能は低下するので、不必要に__syncthreads( )を入れることは望ましくない。

また、Kepler GPUでは同一グループの32スレッドのレジスタファイルの間でデータの交換を行うシャッフル命令が新設された。この命令を使えば、例えば、n=0～31で、スレッドn+1のレジスタiの内容をスレッドnのレジスタiに転送するというスレッド間のレジスタのシフト操作が1命令で実行できる。この演算は整数演算などと同程度の時間で実行でき、シェアードメモリ経由のスレッド間通信より高速である。図3-21にシャッフル命令の幾つかの例を示すが、シャッフル命令ではこれ以外にも多くの転送パターンがサポートされている。

図3-21 KeplerのShuffle命令の一部の転送パターン。本当は32レジスタであるが、紙面の都合で8レジスタの図になっている

キャッシュの必要性

しかし、GPUが汎用の科学技術計算に使われるようになると、大きな2次元や3次元のデータを扱い、同一SMで動かすスレッドグループの範囲を超えてデータアクセスを行うことが必要になってくる。

これをデバイスメモリから少量のデータを読み出してシェアードメモリに書き込み、演算処理を行って、結果の少量のデータをシェアードメモリからデバイスメモリに書き戻すという風にプログラムすることは可能であるが、毎回の転送を記述するのは煩雑であり、プログラムも読みにくくなってしまう。このため、2009年11月に発表されたNVIDIAのFermi GPUではSM単位に1次データキャッシュを設けた。また、グラフィック処理では32bitの単精度浮動小数点演算で十分と言われており、それ以前のGT200 GPUは倍精度浮動小数点演算の性能は単精度の1/8とかなり低い性能であったが、Fermi GPUは単精度の半分の倍精度浮動小数点演算性能を持っている。

さらに、デバイスメモリだけでなく、シェアードメモリ、1次キャッシュ、レジスタファイルなどにECCを付け1bitエラーの訂正と2bitエラーの検出を可能とするなど、Fermi GPUは、世界初の本格的な大規模科学技術計算用のGPUとなった。

図3-22に示すように、シェアードメモリは、メモ用紙のようなもので、手近に置いて、メモしたものを読んだり、メモを書き込んだりすることはできるが、転送が終わると元のデータとの関係は切れてしまう。一方、キャッシュの場合は、元のデータとの対応性が維持されていて、コピーした内容が変更されても元のデバイスメモリに反映されるという違いがある。

また、キャッシュへのデータの出し入れはハードウェアが自動的にやってくれるが、シェアードメモリの場合は、デバイスメモリのどのアドレスのデータをシェアードメモリのどのアドレスにコピーするか、シェアードメモリのどのアドレスのデータをデバイスメモリにコピーするかは、プログラムに明示的に書いてやる必要がある。

図3-22 シェアードメモリとキャッシュメモリの違い

必要なデータだけを移動するという点ではシェアードメモリは効率が良いが、転送を明示的に書く必要があるので、プログラマの負担は大きい。しかし、グラフィックス表示だけにGPUが使われていた時代は、GPUのプログラムを作るのはごく限られたプログラマで、プログラムの手間が掛かっても、ハードウェアのコストを上げずに、高い性能を実現することが望まれたのである。