18

CUDA 同期について質問があります。特に、if ステートメントでの同期について明確にする必要があります。つまり、ブロック内のスレッドの一部がヒットする if ステートメントのスコープに __syncthreads() を配置すると、どうなるでしょうか? 一部のスレッドは、同期点に到達しない他のスレッドを「永遠に」待機し続けると考えました。そこで、検査するサンプル コードをいくつか作成して実行しました。

__global__ void kernel(float* vett, int n)
{
    int index = blockIdx.x*blockDim.x + threadIdx.x;
    int gridSize = blockDim.x*gridDim.x;

    while( index < n )
    {   
        vett[index] = 2;
        if(threadIdx.x < 10)
        {
            vett[index] = 100;
            __syncthreads();
        }
        __syncthreads();

        index += gridSize;
    }
}

驚くべきことに、出力がかなり「正常」であることがわかりました (64 要素、ブロックサイズ 32)。

100 100 100 100 100 100 100 100 100 100 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2
100 100 100 100 100 100 100 100 100 100 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2

そこで、次のようにコードを少し変更しました。

__global__ void kernel(float* vett, int n)
{
    int index = blockIdx.x*blockDim.x + threadIdx.x;
    int gridSize = blockDim.x*gridDim.x;

    while( index < n )
    {   
        vett[index] = 2;
        if(threadIdx.x < 10)
        {
            vett[index] = 100;
            __syncthreads();
        }
        __syncthreads();
            vett[index] = 3;
        __syncthreads();

        index += gridSize;
    }
}

出力は次のとおりです。

3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 
3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3

繰り返しますが、私は間違っていました。ベクトルの要素を変更した後、if ステートメント内のスレッドは待機状態のままであり、if スコープから出ることはないと考えていました。それで...何が起こったのか明確にしていただけますか？同期ポイントの後に取得するスレッドは、バリアで待機しているスレッドのブロックを解除しますか? 私の状況を再現する必要がある場合は、CUDA Toolkit 5.0 RC と SDK 4.2 を使用しました。よろしくお願いします。

4

5 に答える 5

19

つまり、動作は未定義です。そのため、必要な処理が実行される場合と実行されない場合があります。または、カーネルがハングしたりクラッシュしたりする可能性があります。

内部で物事がどのように機能するかを本当に知りたい場合は、スレッドが独立して実行されるのではなく、一度にワープ（32スレッドのグループ）が実行されることを覚えておく必要があります。

もちろん、これは、条件がワープ全体で均一に評価されない条件分岐で問題を引き起こします。この問題は、両方のパスを次々に実行することで解決されます。それぞれのパスでは、そのパスを実行するはずのないスレッドが無効になっています。既存のハードウェア上のIIRCが最初に分岐を取得し、次に分岐が取得されない場所でパスが実行されますが、この動作は定義されていないため、保証されません。

この個別のパスの実行は、コンパイラが2つの個別の実行パスのすべてのスレッド（「再収束ポイント」または「同期ポイント」）が到達することが保証されていると判断できるポイントまで続きます。最初のコードパスの実行がこのポイントに達すると、それは停止され、代わりに2番目のコードパスが実行されます。2番目のパスが同期ポイントに到達すると、すべてのスレッドが再び有効になり、そこから均一に実行が続行されます。

同期の前に別の条件付きブランチが検出されると、状況はさらに複雑になります。この問題は、まだ実行する必要のあるパスのスタックで解決されます（幸い、1つのワープに対して最大32の異なるコードパスを持つことができるため、スタックの成長は制限されます）。

同期ポイントが挿入される場所は未定義であり、アーキテクチャ間でわずかに異なるため、保証はありません。Nvidiaから得られる唯一の（非公式の）コメントは、コンパイラが最適な同期ポイントを見つけるのに非常に優れているということです。ただし、特にスレッドが早期に終了する場合は、最適なポイントが予想よりもさらに下に移動する可能性のある微妙な問題がしばしばあります。

__syncthreads（）ディレクティブ（bar.syncPTXの命令に変換されます）の動作を理解するには、この命令がスレッドごとに実行されるのではなく、ワープ全体に対して一度に実行されることを理解することが重要です（スレッドが無効になっているかどうかに関係なく）ブロックのワープのみを同期する必要があるためです。ワープのスレッドはすでに同期して実行されており、それ以上の同期は効果がないか（すべてのスレッドが有効になっている場合）、異なる条件付きコードパスからスレッドを同期しようとするとデッドロックが発生します。

この説明から、特定のコードの動作に至るまで、さまざまな方法で作業できます。ただし、これはすべて未定義であり、保証はなく、特定の動作に依存すると、いつでもコードが破損する可能性があることに注意してください。

特にコンパイル先の命令については、PTXのマニュアルを参照してください。以下にahmadが参照している、 HenryWongの「マイクロベンチマークによるGPUマイクロアーキテクチャの謎を解き明かす」論文も読む価値があります。現在は時代遅れのアーキテクチャとCUDAバージョンですが、条件分岐に関するセクションは一般的に有効であるように見えます。bar.sync__syncthreads()__syncthreads()

于 2012-09-20T20:28:03.363 に答える
5

CUDA モデルは MIMD ですが、現在の NVIDIA GPU__syncthreads()はスレッドではなくワープの粒度で実装されています。つまり、これらはwarps inside a thread-block必ずしも同期しているとは限りませんthreads inside a thread-block。__syncthreds()スレッド ブロックのすべての「ワープ」がバリアに到達するか、プログラムを終了するまで待機します。詳細については、Henry Wong の Demistifying ペーパーを参照してください。

于 2012-09-20T20:27:17.747 に答える
3

__syncthreads()常に、1 つのスレッド ブロック内のすべてのスレッドでステートメントに到達しない限り、使用しないでください。プログラミング ガイド(B.6)から:

__syncthreads()条件コードで許可されますが、条件がスレッド ブロック全体で同じように評価される場合のみです。そうしないと、コードの実行がハングしたり、意図しない副作用が発生したりする可能性があります。

基本的に、コードは整形式の CUDA プログラムではありません。

于 2012-09-20T20:07:50.060 に答える
1

__syncthreads() は、ブロック内のスレッドを同期するために使用されます。つまり、ブロック内のすべてのスレッドは、すべてが完了するまで待機してから先に進みます。

ブロック内に if ステートメントに入るスレッドとそうでないスレッドがある場合を考えてみましょう。待機中のスレッドはブロックされます。いつまでも待っています。

一般に、if 条件ステートメントに synchronize を入れるのは適切なスタイルではありません。それを避けるのが最善であり、コードがある場合はコードを再設計してください。同期の目的は、すべてのスレッドが一緒に進行することを確認することですが、そもそも if ステートメントを使用してそれらを除外するのはなぜですか?

ブロック間で同期が必要な場合に追加します。カーネルの再起動が必要です。

于 2012-09-21T07:28:33.900 に答える
0

__syncthreads()if 条件では避けたほうがよいでしょう。for ループと for ループの後にコードを書き直すことができ__syncthreads()ます。

于 2020-09-16T14:11:15.087 に答える