Direct ComputeとCUDAの同期命令

メモリバリアやフェンス命令について、調べて分かったことのまとめです。ドキュメントの理解が怪しいうえに、実装して実験していないので、間違っている可能性あり

【そもそもバリアが必要なわけ】

データハザード(RAWハザード・WARハザード・WAWハザード)が起きるから。

SIMT型の(≒GPGPUの)並列処理においてバリアがない場合何が起きるか考える。あるスレッドBが別のスレッドAによって書き込まれた値を参照したいシナリオがあったとする。しかし、スレッドAが値の書き込みが終わっているか否かはスレッドBには分からない。したがって、スレッドAによって更新される前の値を参照してしまう可能性がある。更新を確実にするには、スレッドを跨いだ値の参照が起こるよりも前の時点で、同期待ち(バリア)を行えばよい。

【CUDAの場合】

■__syncthread

この命令が実行されるより前のグローバルメモリとデバイスメモリに対するアクセスが完了し(=このスレッドのメモリ書き込みの結果を同じブロック内の他のスレッドが読み出せる状態になり)、かつ同じブロック内の全スレッドがこの命令にたどり着くまで、これ以降の命令について同じブロック内の全スレッドを実行しない

■__threadfence

この命令が実行されるより前のグローバルメモリとシェアードメモリに対するアクセスが完了するまで、これ以降の命令について同じブロック内の全スレッドを実行しない

※グローバルメモリアクセスは(自分のスレッドのメモリ書き込みが)全スレッドに反映されるまで待機、シェアードメモリアクセスは(自分のスレッドのメモリ書き込みが)同じブロック内のスレッドに反映されるまで待機し、両方の待機が解けたら以降の命令を実行する

■__threadfence_block

この命令が実行されるより前のグローバルメモリとシェアードメモリに対するアクセスが完了するまで、これ以降の命令について同じブロック内の全スレッドを実行しない

※グローバルメモリアクセスは同じブロック内のスレッドに反映されるまで待機、シェアードメモリアクセスも同じブロック内のスレッドに反映されるまで待機し、両方の待機が解けたら以降の命令を実行する

■__threadfence_system

この命令が実行されるより前のグローバルメモリ・シェアードメモリ・ページロックメモリに対するアクセスが完了するまで、これ以降の命令について同じブロック内の全スレッドを実行しない

※2GPU間での同期などに利用する

【Direct Computeの場合】

■GroupMemoryBarrier()

この命令が実行されるより前のgroupsharedメモリに対するアクセスが完了するまで、これ以降の命令について同じグループ内の全スレッドを実行しない

※groupsharedメモリのアクセスがなければ、この命令は無意味

■GroupMemoryBarrierWithGroupSync()

この命令が実行されるより前のgroupsharedメモリに対するアクセスが完了し、かつ同じグループ内の全スレッドがこの命令にたどり着くまで、これ以降の命令について同じグループ内の全スレッドを実行しない

■DeviceMemoryBarrier()

この命令が実行されるより前のデバイスメモリ(おそらくUAVなどのBufferのこと)に対するアクセスが完了するまで、これ以降の命令について同じグループ内の全スレッドを実行しない

※デバイスメモリのアクセスがなければ、この命令は無意味

■DeviceMemoryBarrierWithGroupSync()

この命令が実行されるより前のデバイスメモリに対するアクセスが完了し、かつ同じグループ内の全スレッドがこの命令にたどり着くまで、これ以降の命令について同じグループ内の全スレッドを実行しない

■AllMemoryBarrier()

この命令が実行されるより前の全種類のメモリに対するアクセスが完了するまで、これ以降の命令について同じグループ内の全スレッドを実行しない

※メモリアクセスがなければ、この命令は無意味

■AllMemoryBarrierWithGroupSync()

この命令が実行されるより前の全種類のメモリに対するアクセスが完了し、かつ同じグループ内の全スレッドがこの命令にたどり着くまで、これ以降の命令について同じグループ内の全スレッドを実行しない

【推測というか考察というか】

1. __syncthreads()や**MemoryBarrierWithGroupSync()はブロック内の全スレッドがその命令の位置で同期するが、__threadfence()や**MemoryBarrier()は自分のスレッドのみ待機する(GPUはWarp(Wavefront)を最小粒度として並列実行するから、厳密には「自分のスレッドのみ」ではなく「自分のWarp/Wavefrontに属する全スレッドが」待機する)。

2. __syncthreads()はAllMemoryBarrierWithGroupSync()と同じ。ただし、共有メモリの同期についてのみ考えるならば、__syncthreads()はGroupMemoryBarrierWithGroupSync()と同じ。

3. 分岐により一部のスレッドしか__syncthreads()を呼び出さなかったとき、ブロック内の全スレッドが同期するかどうかはわからない。多分同期されないと思う。

【__threadfenceの使い道】

CUDA C Programming Guideに例が記載されていた。他のブロック間でのデータの受け渡しには、__threadfence(あるいはカーネル関数の再立ち上げ)が必要。

【ソース】

MSDN

CUDA C Programming Guide

Microsoft Direct Compute on the 2nd Generation Intel Core processor(PDF)

CUDA II: additional topics(PDF)

広告

コメントを残す

以下に詳細を記入するか、アイコンをクリックしてログインしてください。

WordPress.com ロゴ

WordPress.com アカウントを使ってコメントしています。 ログアウト / 変更 )

Twitter 画像

Twitter アカウントを使ってコメントしています。 ログアウト / 変更 )

Facebook の写真

Facebook アカウントを使ってコメントしています。 ログアウト / 変更 )

Google+ フォト

Google+ アカウントを使ってコメントしています。 ログアウト / 変更 )

%s と連携中