3

もちろん、ifandswitchステートメントによるワープ発散は、GPU では何としてでも回避する必要があります。

しかし、ワープ ダイバージェンス (特定の行を実行するために一部のスレッドのみをスケジュールする) と追加の無駄な演算のオーバーヘッドはどうなるでしょうか?

次のダミーの例を検討してください。

バージョン 1:

__device__ int get_D (int A, int B, int C)
{
    //The value A is potentially different for every thread.

    int D = 0;

    if (A < 10)
        D = A*6;
    else if (A < 17)
        D = A*6 + B*2;
    else if (A < 26)
        D = A*6 + B*2 + C; 
    else 
        D = A*6 + B*2 + C*3;

    return D;
}

対。

バージョン 2:

__device__ int get_D (int A, int B, int C)
{
    //The value A is potentially different for every thread.

    return  A*6 + (A >= 10)*(B*2) + (A < 26)*C + (A >= 26)*(C*3);
}

私の実際のシナリオはより複雑 (より多くの条件) ですが、考え方は同じです。

質問:

バージョン 1) がバージョン 2 より遅いほど、ワープ ダイバージェンスの (スケジューリングにおける) オーバーヘッドが大きいですか?

バージョン 2 はバージョン 1 よりも多くの ALU を必要とし、これらのほとんどは「0 による乗算」で浪費されます (選択された少数の条件のみが 0 ではなく 1 に評価されます)。これは、無駄な操作で貴重な ALU を拘束し、他のワープの命令を遅らせますか?

4

1 に答える 1