0

次のコードでワープ発散のペナルティが何であるかを理解しようとしています。原則として、ワープ発散がどのように処理されるかを理解しています (小さなブランチの述語命令、大きなブランチのワープ投票と分岐 - すべてのワープが一致する場合、それ以外の場合は述語命令と分岐なし、小さなブランチと同じ)。ただし、詳細はわかりません-中断/継続を伴うwhileループがどのように処理されるか。

以下の例では、scrapeEverythingCondition() がレーン X で true と評価されると、次のいずれかが発生します。

  1. 評価は n == N_N になるまで内側のループで進行し、レーン X は常に nops を評価し、i がインクリメントされ、すべてのレーンが連携するようになります。
  2. レーン X を除くすべての人が someMoreWork() を実行し、レーン X は nops を評価し、レーン X は someCostlyInitialization() を実行し、その後 n=0 を実行しますが、他のすべての人は nops を評価します。
  3. 私が考えていなかった他の何か。

コード：

__global__ void chainKernel() {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    while (i < N_I) {
        someCostlyInitialization();
        for(int n = 0; n < N_N; ++n) {
            someStatisticsComputations(n);
            if (scrapEverythingCondition(n)) {
                // Everything we did for current i is no good. Scrap and begin again
                i -= BLOCKS*THREADS;
                break;
            }
            someMoreWork();
        }
        i += BLOCKS*THREADS;
    }
}

PTX にコンパイルして結果のコードを調べてみましたが、複雑すぎました :(

編集：答えてくれたMakuに感謝します。また、コードの周りに散りばめられた昔ながらの printf() を使用して、答えを確認することもできました。どのスレッドがどこで、どのような順序で取得されたかを確認できました。実際、オプション 1 が正しいものです (レーン X は、内側の for ループが使い果たされるまで中断されます)。

4

2 に答える 2

0

この問題に関する興味深いドキュメントを見つけました: pdf

私が理解していることから、制御フローステートメント (を含むbreak) はスレッドの同期ポイントを定義するということです。あなたの場合、 i += BLOCKS*THREADS; レーンXはforループを離れ、他のスレッドが上記の行に到達するのを待ちます。

于 2013-07-02T09:02:26.217 に答える