3

配列サイズ n を単純に削減しようとしているとしましょう。たとえば、1 つの作業単位内に保持されているとします...すべての要素を追加するとします。一般的な戦略は、各 GPU で多数の作業項目を生成して、ツリー内の項目を減らすことのようです。単純に、これは log n のステップを踏むように見えますが、スレッドの最初の波がこれらのスレッドすべてを一発で処理するわけではありませんね。彼らはワープでスケジュールされます。

for(int offset = get_local_size(0) / 2;
      offset > 0;
      offset >>= 1) {
     if (local_index < offset) {
       float other = scratch[local_index + offset];
       float mine = scratch[local_index];
       scratch[local_index] = (mine < other) ? mine : other;
     }
     barrier(CLK_LOCAL_MEM_FENCE);
   }

したがって、32 個のアイテムが並行して追加され、そのスレッドはバリアで待機します。さらに 32 人が行き、バリアで待機します。さらに 32 回進み、すべてのスレッドがツリーの最上位レベルに到達するために必要な n/2 の追加を行うまでバリアで待機し、ループを回避します。涼しい。

これは良さそうに見えますが、おそらく複雑でしょうか? 命令レベルの並列処理が重要であることは理解しているので、1 つのスレッドを生成して次のようなことをしてみませんか

while(i<array size){
    scratch[0] += scratch[i+16]
    scratch[1] += scratch[i+17]
    scratch[2] += scratch[i+17]
    ...
    i+=16
}
...
int accum = 0;
accum += scratch[0]
accum += scratch[1]
accum += scratch[2]
accum += scratch[3]
...

すべての追加がワープ内で行われるようにします。これで、GPU を好きなだけビジー状態に保つスレッドが 1 つになりました。

ここで、命令レベルの並列処理は実際には問題ではないと仮定します。以下、ワークサイズを32（ワープ数）に設定した場合はどうでしょうか。

for(int i = get_local_id(0);i += 32;i++){
    scratch[get_local_id(0)] += scratch[i+get_local_id(0)]
}

次に、最初の 32 項目を合計します。これらの 32 個のスレッドが何度も起動し続けると想像します。

OpenCL の汎用性を放棄することに反対しないのであれば、1 サイクルあたりに実行される add の数がわかっているのに、わざわざツリーを削減する必要があるでしょうか?

4

1 に答える 1