5

ブロック内のすべてのスレッドがコード内の同じポイントにあることが絶対に必要な場合、起動されるスレッドの数がワープ内のスレッドの数と等しい場合、__ syncthreads関数が必要ですか？

注：余分なスレッドやブロックはなく、カーネルのワープは1つだけです。

コード例：

shared _voltatile_ sdata[16];

int index = some_number_between_0_and_15;
sdata[tid] = some_number;
output[tid] = x ^ y ^ z ^ sdata[index];
4

2 に答える 2

8

volatile の使用に関する詳細情報を更新

おそらく、すべてのスレッドが他のスレッドによって共有メモリに書き込まれたデータを読み取っているため、すべてのスレッドが同じポイントにあることを望んでいます。単一のワープを (各ブロックで) 起動している場合、すべてのスレッドが一緒に実行されていることがわかります。__syncthreads()一見すると、これは、「ワープ同期プログラミング」として知られるプラクティスを省略できることを意味します。ただし、注意すべき点がいくつかあります。

  • コンパイラは、データをレジスタに保持できるメモリへの格納を遅らせるなど、スレッド内セマンティクスが正しいままであれば最適化できると想定することに注意してください。__syncthreads()これに対する障壁として機能するため、他のスレッドがデータを読み取る前に、データが共有メモリに書き込まれることが保証されます。を使用するvolatileと、コンパイラはレジスタに保持するのではなくメモリ書き込みを実行しますが、これにはいくつかのリスクがあり、よりハックです (つまり、これが将来どのように影響を受けるかはわかりません)。
    • __syncthreads()技術的には、CUDA プログラミング モデルに準拠するために常に使用する必要があります。
  • ワープ サイズは常に 32 ですが、次のことができます。
    • コンパイル時warpSizeに、デバイス コードで特殊変数を使用します ( CUDA プログラミング ガイドの「組み込み変数」の下、4.1 バージョンのセクション B.4 に記載されています)。
    • 実行時に cudaDeviceProp 構造体の warpSize フィールドを使用します ( CUDA リファレンス マニュアルに記載されています) 。

一部の SDK サンプル (特にリダクションとスキャン) では、このワープ同期手法が使用されていることに注意してください。

于 2012-04-18T10:00:52.677 に答える
1

__syncthreads()ワープが並行して実行されている場合でも必要です。SM (ストリーム マルチプロセッサ) 内のコア数は 32 未満になる可能性があるため、ハードウェアでの実際の実行は並列ではない可能性があります。たとえば、GT200 アーキテクチャには各 SM に 8 つのコアがあるため、すべてのスレッドがコードの同じポイント。

于 2012-04-18T09:15:55.870 に答える