17

どのような状況volatileで、CUDA カーネルの共有メモリでキーワードを使用する必要がありますか? コンパイラに値をキャッシュしないように指示していることは理解していvolatileますが、私の質問は共有配列での動作に関するものです。

__shared__ float products[THREADS_PER_ACTION];

// some computation
products[threadIdx.x] = localSum;

// wait for everyone to finish their computation
__syncthreads();

// then a (basic, ugly) reduction:
if (threadIdx.x == 0) {
    float globalSum = 0.0f;
    for (i = 0; i < THREADS_PER_ACTION; i++)
        globalSum += products[i];
}

productsこの場合、揮発性である必要がありますか? 各配列エントリは、すべてがスレッド 0 によって読み取られる最後を除いて、単一のスレッドによってのみアクセスされます。コンパイラが配列全体をキャッシュできる可能性はあります volatileか? ?

ありがとう！

4

2 に答える 2

23

共有配列を として宣言しない場合volatile、コンパイラは共有メモリ内の場所をレジスタ (スコープは単一のスレッドに固有) に配置することで自由に最適化できます。これは、特定の共有要素に 1 つのスレッドからのみアクセスするかどうかに関係なく当てはまります。したがって、共有メモリをブロックのスレッド間の通信手段として使用する場合は、それを宣言することをお勧めしますvolatile。ただし、この種の通信パターンでは、多くの場合、読み取り/書き込みの順序を強制する実行バリアも必要になるため、以下のバリアについて読み続けてください。

明らかに、各スレッドが共有メモリの独自の要素にのみアクセスし、別のスレッドに関連付けられた要素にはアクセスしない場合、これは問題ではなく、コンパイラの最適化によって何も壊れることはありません。

あなたの場合、各スレッドが共有メモリの独自の要素にアクセスしているコードのセクションがあり、スレッド間アクセスのみがよく理解されている場所で発生する場合、メモリフェンス関数 を使用してコンパイラに強制的に強制的に削除させることができますレジスタに一時的に格納されている値は、共有配列に戻されます。あなたはそれが役に立つかもしれないと思う__threadfence_block()かもしれませんが、あなたの場合、 には__syncthreads() すでにメモリフェンシング機能が組み込まれています。したがって、__syncthreads()スレッドの同期を強制するだけでなく、共有メモリ内のレジスタにキャッシュされた値を強制的に共有メモリに強制的に戻すには、この呼び出しで十分です。

ところで、コードの最後のリダクションがパフォーマンスの問題である場合は、並列リダクション メソッドを使用して高速化することを検討できます。

于 2013-03-11T04:20:57.167 に答える
-1

ここに来る他の人のために簡単に言えば：

呼び出し__syncthreads()は、共有メモリを として宣言するよりも強力volatileです。__syncthreads()特定のワークグループのすべてのスレッドが 1 つの共通点で一緒に停止し、メモリを同期させます。

volatileOTOH は、コンパイラがキャッシュの最適化を実行しないようにすることで (コストがかかる場合があります)、スレッド間で特定のメモリ バッファーの一貫性を保ちますが、各スレッドは独自のペースで自由に移動できるため、コンパイラ/ハードウェアはあらゆる種類のスケジューリングを実行できます。最適化。
(ただし、書き込みが複数のプロセッサ命令から構成される場合、volatile はデータの整合性を保証しないことに注意してください)

要約すると、必要なのはスレッド間のメモリの一貫性だけで、すべてを 1 点で停止しない場合は、volatile通常、よりも優れたパフォーマンスが得られ__syncthreads()ます。ただし、特定のアルゴリズムや入力データに応じてミレージが異なる場合があるため、パフォーマンスの最後のビットをすべて絞り込む必要がある場合は、両方のアプローチをテストしてください。

さらに、ワークグループ内のアクティブなスレッドの数が SIMD 幅 (ワープサイズ) より小さい場合、同じワープ内のすべてのスレッドが同期的に命令を実行するvolatile代わりに、 を使用できます。__synchthreads()たとえば、最後のラップ アンローリングの最適化を並列削減アルゴリズム(スライド 21 ～ 23) に参照してください。これは、アクティブなスレッドの数がワープ サイズよりも小さくなった場合に__synchthreads()のみ最初に使用し、後で依存します。volatile

于 2021-10-08T06:18:13.990 に答える