単一の Warp のスレッド間のアトミックな競合 (並行性) と、1 つのブロック内の異なる Warp のスレッド間のアトミックの競合 (並行性) のどちらが優れていますか? 共有メモリにアクセスするときは、別のワープのスレッドよりも、1 つのワープのスレッド同士の競合が少ない方がよいと思います。逆にグローバルメモリにアクセスする場合は、1 つのブロックの異なるワープのスレッドが、単一のワープのスレッドよりも競合が少ない方がよいのではないでしょうか。
競合 (同時実行性) を解決する方法と、単一のワープ内のスレッド間またはワープ間でストアを分離する方法を知る必要があります。
ちなみにチームは __ syncthreads (); と言えるかもしれません。1 つのワープのスレッドではなく、単一のブロックでワープを同期しますか?
ブロック内のかなりの数のスレッドが同じ値に対してアトミック更新を実行する場合、これらのスレッドはすべてシリアル化する必要があるため、パフォーマンスが低下します。このような場合、通常は、各スレッドがその結果を別の場所に書き込み、別のカーネルでそれらの値を処理する方が適切です。
ワープ内の各スレッドが同じ値へのアトミック更新を実行する場合、ワープ内のすべてのスレッドが同じクロック サイクルで更新を実行するため、アトミック更新の時点ですべてのスレッドをシリアル化する必要があります。これはおそらく、ワープがすべてのスレッドを処理するために 32 回スケジュールされていることを意味します (非常に悪い)。
一方、ブロック内の各ワープの 1 つのスレッドが同じ値へのアトミック更新を実行する場合、ワープのペア (2 つのワープ スケジューラによって各クロックで処理される 2 つのワープ) がオフセットされるため、影響は小さくなります。処理パイプラインを移動する際に (1 クロック サイクルずつ)。したがって、最終的には 2 つのアトミック更新 (2 つのワープのそれぞれから 1 つ) だけになり、1 サイクル内で発行され、すぐにシリアル化する必要があります。
したがって、2 番目のケースでは、状況は改善されますが、それでも問題があります。その理由は、共有値がどこにあるかによっては、SM 間でシリアル化を行うことができますが、これは非常に遅くなる可能性があるためです。これは、各スレッドが更新がグローバル メモリに送信されるまで待機する必要がある場合があるためです。 L2、そして戻る。ブロック内のスレッドが共有メモリ (L1) 内の値に対してアトミック更新を実行し、次に各ブロック内の 1 つのスレッドがグローバル メモリ (L2) 内の値に対してアトミック更新を実行するように、アルゴリズムをリファクタリングできる場合があります。 ）。
アトミック操作は完全な命の恩人になる可能性がありますが、CUDA を初めて使用する人によって過度に使用される傾向があります。多くの場合、並列リダクションまたは並列ストリーム圧縮アルゴリズムで別のステップを使用することをお勧めします (「参考文献」を参照thrust::copy_if)。