3

大きな配列の削減は、__ reduce（）;を呼び出すことで実行できます。複数回。

ただし、次のコードは2つのステージのみを使用し、ここに文書化されています。

ただし、この2段階の削減のアルゴリズムを理解できません。もっと簡単な説明ができる人もいますか？

__kernel
void reduce(__global float* buffer,
        __local float* scratch,
        __const int length,
        __global float* result) {

    int global_index = get_global_id(0);
    float accumulator = INFINITY;
    // Loop sequentially over chunks of input vector
    while (global_index < length) {
        float element = buffer[global_index];
        accumulator = (accumulator < element) ? accumulator : element;
        global_index += get_global_size(0);
    }

    // Perform parallel reduction
    int local_index = get_local_id(0);
    scratch[local_index] = accumulator;
    barrier(CLK_LOCAL_MEM_FENCE);
    for(int offset = get_local_size(0) / 2; offset > 0; offset = offset / 2) {
        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);
    }
    if (local_index == 0) {
        result[get_group_id(0)] = scratch[0];
    }
}

CUDAを使用して適切に実装することもできます。

4

1 に答える 1

5

Nスレッドを作成します。最初のスレッドは位置0、N、2 * N、...の値を調べます。2番目のスレッドは値1、N + 1、2 * N + 1、...を調べます。これが最初のループです。length値をN個の値に減らします。

次に、各スレッドは最小値を共有/ローカルメモリに保存します。次に、同期命令（barrier(CLK_LOCAL_MEM_FENCE)。）があります。次に、共有/ローカルメモリの標準的な削減があります。ローカルID0のスレッドが完了すると、その結果が出力配列に保存されます。

length全体として、からの値への削減がありN/get_local_size(0)ます。このコードの実行が完了したら、最後のパスを1回実行する必要があります。ただし、これでほとんどの作業が完了します。たとえば、長さが10 ^ 8、N = 2 ^ 16、get_local_size（0）= 256 = 2 ^ 8の場合、このコードは10^8要素を256要素に減らします。 。

わからない部分は？

于 2012-06-11T07:28:41.540 に答える