CUDAスレッド同期

CUDAスレッドについて†

カーネル呼び出し(<<< >>>)で複数のスレッドを同時に立ち上げたとき， これらのスレッドは同時には実行されません．スレッドの呼び出しには1クロック 必要なので，例えば，100スレッドを一斉に動かすと，最初と最後のスレッドでは 100クロックずれます． このような非同期のスレッドを同期させるのが__syncthreads()であり， また，メモリアクセスに対する排他制御を行うのが__threadfence()，__threadfence_block()です．

__syncthreads()†

CUDAで行列演算：乗算(シェアードメモリ使用版)でも説明していますが， カーネル関数中に__syncthreads()を記述すると， その行が同期点として設定されます． 各スレッドは処理が同期点に達すると，他のスレッドがその同期点に達するまで待機します． 同期点がバリアのような働きをするのでバリア同期と呼ばれます． 注意点として，__syncthreads()によるバリアは同一ブロック内のスレッドのみ有効です． これは，性能の異なる複数の異なるデバイスを組み合わせて実行した場合に， 処理の速いデバイスが処理の遅いデバイスに合わせられ， 全体のパフォーマンスが低下することを防ぐためです． 異なるブロックの同期(グローバル同期)が必要な場合は，一度カーネルを終了させることで対応します．

__threadfence()，__threadfence_block()†

void __threadfence();

は__threadfence()より前に呼ばれたスレッドによるメモリアクセス(グローバルメモリとシェアードメモリ)全てが デバイス内(シェアードメモリの場合はブロック内)の全てのスレッドで確認できるようになるまで待機します(Programing Guideまま)． データの書き込み中に他のスレッドからのアクセスを防いでくれます．__threadfence()がデバイス内の全スレッド， __threadfence_block()はブロック内の全スレッドが対象です．