グローバルメモリにシーケンシャルにアクセスするスレッドが多数あると仮定すると、どのオプションが全体的に高速に実行されますか? __threadfence() はすべての共有メモリとグローバルメモリの書き込みを考慮に入れていますが、書き込みは結合されているため、私には疑問があります。一方、atomicExch() は重要なメモリ アドレスだけを考慮しますが、書き込みが結合されているかどうかはわかりません。
コード内:
array[threadIdx.x] = value;
または
atomicExch(&array[threadIdx.x] , value);
ありがとう。
atomicExchKepler GPU では、アトミックが Kepler で非常に高速であるため、私は賭けます。Fermi では、ウォッシュかもしれませんが、衝突がないことを考えると、atomicExchそれでもうまく機能する可能性があります。
実験して結果を報告してください。
この 2 つは非常に異なることを行います。
atomicExch一度に 2 つのスレッドが特定のセルを変更しようとしないことを保証します。このような競合が発生すると、1 つ以上のスレッドが停止する可能性があります。2 つのスレッドが同じセルにアクセスしないことが事前にわかっている場合は、atomic...関数を使用する意味がありません。
__threadfence()現在のスレッド (および現在のスレッドのみ!) を遅延させて、特定のスレッドによる後続の書き込みが実際に後で発生するようにします。そのため、__threadfence()フォローアップ コードがなければ、それ自体はあまり興味深いものではありません。
そのため、その 2 つの効率を比較する意味はないと思います。もう少し具体的なユースケースを示すことができれば、私は関連付けることができます...
どちらも実際には、スレッドの実際の実行順序を保証するものではないことに注意してください。