コンピュータアーキテクチャの話 (331) 汎用の科学技術計算におけるGPUのキャッシュの必要性

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

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