0

私のカーネルはとてもシンプルです。コードが有効かどうかを確認し、プレフィックス スキャンの出力に従って一意のコードのみを格納します。

__kernel void moveValid(__global int* sortCode, __global int* mark, __global int* processorOffsets, __global int* uniqueCode,__global int* numPoints, __global int* pointIndex) 
{
    int ig = get_global_id(0);
    int m = mark[ig];
    int j= processorOffsets[ig];
    atomic_inc(&numPoints[j-1]);
    // select
    if(m == true)
    {   
            uniqueCode[j] = sortCode[ig];
            pointIndex[j] = ig;
    }

    barrier(CLK_GLOBAL_MEM_FENCE);
}

カーネルが本当に遅いようです。if文のせいでしょうか？カーネルを改善する方法について誰かヒントを教えてもらえますか? また、このシナリオで select を使用できますか?

4

1 に答える 1

2

したがって、コードを深く調べなくても、その速度について次のフィードバックを提供できます。デバイスとして GPU を使用していると仮定しています。デバイスとして CPU を使用している場合でも、情報の一部が適用される場合があります。

atomic_inc(&numPoints[j-1]);

アトミックインクリメントは、ほとんどのデバイスでグローバルメモリに対して非常に遅くなります。これは、そのデータをグローバル メモリにコミットする必要があるためです (ローカルにキャッシュすることはできません)。

バリア (CLK_GLOBAL_MEM_FENCE);

このバリアは、work_group 内のすべての work_items が実行を継続する前に存在することを保証します。なぜコードでこれが必要なのですか? 特に何もすることが残っていない場合、スレッドの実行を終了させない理由はありません。これも大きなパフォーマンスヒットです。

もし(m ==真)

これは、分岐が 1 つしかないため、実際には私が見た中で最悪の if ステートメントではありません。これでもコードは遅くなりますが、他のものほどではありません。続行する前に、(一部のアーキテクチャでは) すべてのスレッドの条件が順次計算されます。

全体として、このコードでは、アトミック操作と数学操作なしで 4 つのグローバル メモリ アクセスを実行しています。GPU は、このタイプのアルゴリズムを実行するには最悪のタイプのデバイスです。これは、特にアクセスが結合されている場合、メモリ アクセスがグローバル メモリに対して非常に遅いためです。代わりに、配列の一部をローカル メモリに移動することを検討していただけますか?

于 2013-09-13T21:10:49.430 に答える