2

CUDA でプログラムを書き込もうとしていますが、スレッド間の同じブロックでの同期に問題があります。

モデルの状況は次のとおりです。

 10 __global__ void gpu_test_sync()
 11 {
 12     __shared__ int t;
 13     int tid = threadIdx.x;
 14
 15     t = 0;
 16     __threadfence();
 17     __syncthreads();
 18
 19     // for(int i=0; i<1000000 && t<tid; i++); // with fuse
 20     while(t<tid);
 21
 22     t++;
 23     __threadfence();
 24 }
 25
 26 void f_cpu()
 27 {
 28     printf("TEST ... ");
 29     int blocks = 1;
 30     int threads = 2;
 31     gpu_test_sync<<< blocks , threads >>>();
 32     printf("OK\n");
 33 }

スレッド = 1 の場合、すべて問題ありません。スレッド数が 1 より大きい場合、無限循環。

なんで？関数 __threadfence(); 他のスレッドの t 変数の値を可視化する必要があります。

どうすれば解決できますか？

4

3 に答える 3

7

while(t<tid)ワープのすべてのスレッドが無期限にループし、行に到達しない分岐分岐のために、カーネルがあなたがしようとしていることを実行できるとは思いません++t。

長い説明

スレッド、ブロック、ワープについて既に知っている場合は、重要なものについて「重要な部分」までスクロールします。

(私はまだ Kepler アーキテクチャの経験がありません。Fermi を使用していない場合、これらの数値の一部は異なる場合があります。)

次のセクションを理解するために、いくつかの用語について説明する必要があります。

  • スレッド – 単一の実行スレッド。
  • ブロック – 同じカーネルを実行する複数のスレッドのグループ。
  • grid – ブロックのグループ。

次の用語は、物理 (ハードウェア アーキテクチャに依存する物理) スレッドに関連しています。

  • コア – 単一の計算コア。1 つのコアは一度に 1 つの命令を正確に実行します。
  • ワープ – ハードウェア上で並行して実行されるスレッドのグループ。ワープは、現世代の CUDA ハードウェア上の 32 のスレッドで構成されます。

カーネルは、1 つ以上のストリーミング マルチプロセッサ (SM) によって実行されます。Fermi ファミリ (GeForce 400 および GeForce 500 シリーズ) の典型的なミッドエンドからハイエンドの GeForce カードには、単一の GPU に 8 ～ 16 個の SM があります [ Fermi ホワイトペーパー]。各 SM は 32 個の CUDA コア (コア) で構成されます。スレッドは、ワープ スケジューラによって実行されるようにスケジュールされます。各 SM には、ロックステップ方式で動作する 2 つのワープ スケジューラ ユニットがあります。ワープ スケジューラがスケジュールできる最小単位はワープと呼ばれ、執筆時点でこれまでにリリースされたすべての CUDA ハードウェアの 32 スレッドで構成されます。各 SM で一度に実行できるワープは 1 つだけです。

CUDA のスレッドは CPU スレッドよりもはるかに軽量であり、コンテキスト スイッチは安価であり、ワープのすべてのスレッドが同じ命令を実行するか、ワープの他のスレッドが命令を実行する間待機する必要があります。これは、単一命令マルチ スレッド ( SIMT) であり、SSE、AVX、NEON、Altivec などの従来の CPU SIMD (Single Instruction Multiple Data) 命令に似ています。

解決するために 32 を超えるスレッドを必要とする問題を可能にするために、CUDA スレッドは、ソフトウェア開発者によって定義されたサイズのブロックおよびグリッドと呼ばれる論理グループに配置されます。ブロックはスレッドの 3 次元の集合であり、ブロック内の各スレッドには独自の 3 次元識別番号があり、開発者はカーネル コード内のスレッドを区別できます。単一ブロック内のスレッドは、共有メモリを介してデータを共有できます。これにより、グローバル メモリの負荷が軽減されます。共有メモリは、グローバル メモリよりもレイテンシがはるかに低くなりますが、限られたリソースです。ユーザーは、(ブロックごとに) 16 kB の共有メモリと 48 kB の L1 キャッシュ、または 48 kB の共有メモリと 16 kB の L1 キャッシュのいずれかを選択できます。

スレッドのいくつかのブロックは、グリッドにグループ化できます。グリッドは、ブロックの 3 次元配列です。最大ブロック サイズは利用可能なハードウェア リソースに関連付けられていますが、グリッドは (ほぼ) 任意のサイズにすることができます。グリッド内のブロックは、レイテンシが最も高いオン GPU メモリであるグローバル メモリを介してのみデータを共有できます。

Fermi GPU は、SM ごとに一度に 48 のワープ (1536 スレッド) をアクティブにすることができます。これは、スレッドがローカル メモリと共有メモリをほとんど使用せず、同時にすべてに適合するためです。レジスタがスレッドに割り当てられているため、スレッド間のコンテキストの切り替えは高速であり、したがって、スレッドの切り替え間でレジスタと共有メモリを保存および復元する必要はありません。その結果、ストールが発生するたびにワープ スケジューラが現在アクティブなワープを切り替えることにより、カーネル内のメモリ ストールが隠蔽されるため、ハードウェアを過剰に割り当てることが実際に望まれます。

重要な部分

スレッド ワープは、同じストリーミング マルチプロセッサ(SM) 上で実行されるスレッドのハードウェア グループです。ワープのスレッドは、スレッド間で共通のプログラム カウンターを共有することと比較できます。したがって、すべてのスレッドがプログラム コードの同じ行を実行する必要があります。コードに次のような分岐ステートメントがある場合if ... then ... elseワープは最初のブロックに入るスレッドを最初に実行する必要があり、ワープの他のスレッドは待機し、次のブロックに入るスレッドは他のスレッドが待機している間に実行されます。この動作のため、GPU コードでは条件付きステートメントをできるだけ避ける必要があります。ワープのスレッドが異なる実行ラインに従う場合、それは分岐スレッドを持つと呼ばれます。条件付きブロックは CUDA カーネル内で最小限に抑える必要がありますが、同じワープのすべてのスレッドがif ... then ... elseブロック内の単一の実行パスのみをたどるようにステートメントを並べ替えて、この制限を緩和できる場合があります。

whileandforステートメントは分岐ステートメントなので、 に限定されませんif。

于 2012-11-15T12:23:41.637 に答える
2

複数のスレッドでカーネルを起動すると、ゼロより大きいwhile(t<tid);スレッドの無限ループが発生するため、無限ループが発生します。idx

この時点で、問題はスレッドの同期ではなく、実装したループに関連しています。

于 2012-11-15T12:43:03.023 に答える
1

一連のスレッドを連続して実行しようとしている場合は、CUDA を悪用しています。

また、最初のスレッドを過ぎたスレッドは更新された t を受信しないため、機能しません。共有変数を更新するには __syncthreads() を呼び出す必要がありますが、それを行うことができるのは、すべてのスレッドが同じことを実行している場合のみです。 - つまり、待っていません。

于 2012-11-15T12:29:46.793 に答える