1

GPU で 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
        while (global_index < length) {
            float element = buffer[global_index];
            accumulator = (accumulator < element) ? accumulator : element;
            global_index += get_global_size(0);
    }
    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];
    }
}
4

1 に答える 1

0 
accumulator = (accumulator < element) ? accumulator : element;

fmin関数を使用してください- それはまさにあなたが必要とするものであり、より高速なコードになる可能性があります (コストのかかる分岐の代わりに、利用可能な場合は組み込み命令を呼び出します)

global_index += get_global_size(0);

あなたの典型は何get_global_size(0)ですか？

アクセス パターンはそれほど悪くはありませんが (32 ワープでは 128 バイトのチャンクが結合されています)、可能な限りメモリに順次アクセスすることをお勧めします。たとえば、シーケンシャル アクセスはメモリのプリフェッチOpenCLに役立ちます (コードは、CPU を含む任意のデバイスで実行できることに注意してください)。

次のスキームを検討してください。各スレッドは範囲を処理します

[ get_global_id(0)*delta ,  (get_global_id(0)+1)*delta )

これにより、完全に順次アクセスが行われます。

于 2013-04-28T12:14:42.807 に答える