5

更新:while()以下の条件はコンパイラによって最適化されるため、両方のスレッドは単に条件をスキップし、-O0フラグを付けても CS に入ります。コンパイラがこれを行っている理由を誰かが知っていますか? ちなみに、グローバル変数を宣言するvolatileと、奇妙な理由でプログラムがハングアップします...

CUDA プログラミング ガイドを読みましたが、 CUDAがグローバル メモリに関してメモリの一貫性をどのように処理するかについては、まだ少し不明です。(これはメモリ階層とは異なります) 基本的には、順次整合性を崩そうとするテストを実行しています。私が使用しているアルゴリズムは、カーネル関数内の2 つのスレッド間の相互排除のための Peterson のアルゴリズムです。

flag[threadIdx.x] = 1; // both these are global
turn = 1-threadIdx.x;

while(flag[1-threadIdx.x] == 1 && turn == (1- threadIdx.x));
shared_gloabl_variable_x ++;

flag[threadIdx.x] = 0;

これはかなり簡単です。各スレッドは、フラグを 1 に設定し、別のスレッドに順番を与えることで、クリティカル セクションを要求します。の評価時にwhile()、他のスレッドがそのフラグを設定していない場合、要求元のスレッドは安全にクリティカル セクションに入ることができます。このアプローチの微妙な問題は、コンパイラが書き込みの順序を変更して、 への書き込みが へturnの書き込みの前に実行される場合flagです。これが発生すると、両方のスレッドが同時に CS で終了します。ほとんどのプロセッサは順次整合性を実装していないため、通常の Pthreads でこれを証明するのはかなり簡単です。しかし、GPU はどうでしょうか?

これらのスレッドは両方とも同じワープになります。そして、ステートメントをロックステップ モードで実行します。しかし、それらがturn変数に到達すると、同じ変数に書き込まれるため、ワープ内の実行はシリアル化されます (順序は関係ありません)。この時点で、勝ったスレッドは while 条件に進みますか、それとも他のスレッドが書き込みを終了するのを待って、両方がwhile()同時に評価できるようにしますか? パスは で再び分岐しwhile()ます。これは、そのうちの 1 つだけが勝ち、もう 1 つが待機するためです。

コードを実行した後、一貫して SC を壊すようになりました。私が読んだ値は常に 1 です。これは、両方のスレッドが何らかの形で毎回 CS に入っていることを意味します。これはどのように可能ですか (GPU は命令を順番に実行します)? (注: でコンパイルした-O0ため、コンパイラの最適化は行われず、したがって は使用されませんvolatile)。

4

2 に答える 2

3

編集: スレッドが 2 つしかなく1-threadIdx.x動作するため、スレッド ID 0 と 1 を使用する必要があります。スレッド 0 と 1 は、現在のすべての NVIDIA GPU で常に同じワープの一部になります。ワープは、発散条件のスレッド実行マスクを使用して、SIMD 方式で命令を実行します。while ループは発散状態です。

  • turnとflags が でない 場合volatile、コンパイラはおそらく命令を並べ替え、CS に入る両方のスレッドの動作を確認します。
  • turnとflags が の場合volatile、ハングが表示されます。その理由は、スレッドの 1 つがターンの書き込みに成功するためturn、0 または 1 のいずれかになります。仮定しましょうturn==0: ハードウェアが分岐分岐のスレッド 0 の部分を実行することを選択した場合、すべて問題ありません。しかし、分岐分岐のスレッド 1 の部分を実行することを選択した場合、while ループでスピンし、スレッド 0 は決して順番を取得しないため、ハングします。

2 つのスレッドが異なるワープにあることを確認することでハングを回避できる可能性がありますが、両方から命令を発行して進行できるように、SM にワープを同時に常駐させる必要があると思います。(これはグローバル メモリであるため、異なる SM での同時ワープで動作する可能性がありますが、それには __threadfence_block() だけでなく __threadfence() が必要になる場合があります)。

一般に、これは、このようなコードが GPU では安全でなく、使用すべきでない理由の良い例です。ただし、これは単なる調査実験であることは理解しています。一般に、CUDA GPU は、ほとんどのプロセッサが実装していないと述べているように、順次整合性を実装していません。

元の回答

  1. そうしないと、 のロードが繰り返されず、条件turnが再評価されず、代わりに と見なされます。flagvolatileflagturn == 1-threadIdx.Xtrue
  2. 正しい順序を得るには、store toと store toの__threadfence_block()間にが必要です。flagturn
  3. 共有変数のインクリメントの前に a が必要__threadfence_block()です (これも宣言する必要がありますvolatile)。他のスレッドから確実に見えるようにするために、 __syncthreads()少なくともインクリメントの後に または が必要な場合もあります。__threadfence_block()

ただし、これらの修正を行った後でも、まだ問題が発生する可能性があるという予感があります。それがどうなるか教えてください。

ところで、この行には構文エラーがあるため、これが実際のコードではないことは明らかです。

while(flag[1-threadIdx.x] == 1 and turn==[1- threadIdx.x]);
于 2012-04-23T04:26:44.393 に答える
2

__threadfence() などの余分なメモリ バリアがない場合、グローバル メモリの順次整合性は特定のスレッド内でのみ適用されます。

于 2012-04-22T17:58:47.743 に答える