CUDAスレッド同期の変更点

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

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

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