2

最近、CUDA内のブロックのスレッド間の同期の問題について質問しました。ここで、スレッドを早期に終了すると、ブロック内のCUDAスレッド間の同期が中断されますか?私の質問へのコメントの1つは、同様のスレッドへのリンクを提供しました。これは、PTXガイドのCUDAバリア(__syncthreads())命令について次のように引用しています。

バリアは、ワープ内のすべてのスレッドがアクティブであるかのように、ワープごとに実行されます。したがって、ワープ内のいずれかのスレッドがバー命令を実行した場合、ワープ内のすべてのスレッドがバー命令を実行したかのようになります。ワープ内のすべてのスレッドは、バリアが完了するまでストールし、バリアの到着カウントは、ワープサイズ(ワープ内のアクティブなスレッドの数ではない)だけ増加します。条件付きで実行されるコードでは、bar命令は、すべてのスレッドが条件を同じように評価することがわかっている場合にのみ使用する必要があります(ワープは発散しません)。バリアはワープごとに実行されるため、オプションのスレッド数はワープサイズの倍数である必要があります。

この引用で説明されているメカニズムについては、まだ少し混乱しています。条件付きコードでバリアを使用し、条件付きコードで別のパスを使用して一部のスレッドがバリア命令に到達できない場合、未定義の動作やデッドロックが発生する可能性があることを示しています。問題は、このメカニズムがどのようにデッドロックを引き起こす可能性があるのか​​理解していないということです。(ワープサイズの倍数ではないスレッド番号でも危険です。)このドキュメントでは、1つのスレッドでもバー命令を実行すると、ワープ内のすべてのスレッドがワープ命令と到着を実行したかのように扱われると記載されています。カウンタは、ワープのスレッド数によって更新されます。おそらく、CUDAアーキテクチャは、この到着カウンターをチェックすることによって、すべてのスレッドが同期されているかどうかを判断します。ブロック内の実際のスレッド数と比較します。スレッドごとに更新された場合、カウンターが最大に達することはないため、デッドロックが発生する可能性があります。num。それらのいくつかはbar命令を含まない条件付きパスをとったので、スレッドの数。ただし、ここでは、ワープのスレッド数で数が更新されます。ですから、私はここで根本的なメカニズムを正確に理解していません。

私の他の質問は、全体的な条件文についてです。ワープ内のすべてのスレッドが特定の時間に同じ命令を実行することを知っています。if句の場合、if分岐とelse分岐をとるスレッドは、アイドル状態を維持し、条件の最後に再度同期することで、互いに待機します。したがって、このような条件付きコードには暗黙の同期メカニズムがあります。さて、これは次のようなコードでどのように機能しますか?

int foundCount=0;
for(int i=0;i<array1_length;i++)
{
    for(j=0;j<array0_length;j++)
    {
        if(i == array0[j])
        {
            array1[i] = array1[i] + 1;
            foundCount++;
            break;
        }
    }

     if(foundCount == foundLimit)
        break;
}

これは私の現在のタスクからのコードの一部です。array1の各メンバーについて、現在のarray1インデックスがarray0に含まれているかどうかを確認する必要があります。もしそうなら、私はarray1の現在のインデックスの要素をインクリメントし、それはすでにarray0に含まれているので、breakステートメントで内部ループを終了します。array1に含まれるインデックスの総数が制限に達した場合、外側のループを続行する必要はなく、それを終了することもできます。これはCPUコードでは簡単ですが、CUDAのワープメカニズムがこのようなネストされた条件付きケースをどのように処理するかを知りたいです。32スレッドのワープがこのコードを処理していると想像してください。一部のスレッドは内部ループを処理できますが、一部はすでに終了しており、一部は外部ループからも終了している可能性があります。この場合、アーキテクチャはスレッドの動作をどのように編成しますか?スレッドの現在の「待機ポイント」のリストを保持していますか?このような複雑な状況で、同じワープのスレッドが同じコード行を処理することをどのように保証しますか?

4

1 に答える 1

3

条件分岐は、ワープ内のすべてのスレッドがすべての分岐を実行することによって実装されます。分岐に従わないスレッドは、null op と同等の処理を実行します。これは通常、マスクされた実行と呼ばれ、部分的なワープに対応する方法でもあります。部分的なワープには、永続的にマスクされたスレッドが含まれます。分岐なしで三項演算子などを実装するために利用できる直接条件付き実行命令もあります。

barこれらのメカニズムは、標準のPTX 命令には適用されません。お気づきのように、これは単純なカウンター デクリメント スキームを使用して実装されており、ブロック内のすべてのスレッドがカウンターをゼロにデクリメントしない場合、デッドロックが発生します。

于 2012-07-27T12:41:41.503 に答える