search - CUDA での初出検索

Question

私のアプリケーションは でいくつかの処理を行いdevice-code、 内に配列を生成しkernelます。

この配列内で最初に出現する要素を検索する必要があります。GPUでどのように実行できますか? 配列を CPU にコピーしてそこで作業を行うと、このコードが何度も呼び出されるため、非常に多くのメモリ トラフィックが生成されます。

Answer 1

おそらくもっと洗練された解決策がありますが、最初は特に要素の出現回数が非常に少ない場合は、単純な力ずくのatomic-minが実行可能な解決策になる可能性があります。

template<typename T> __global__ void find(T *data, T value, int *min_idx)
{
    int idx = threadIdx.x + blockDim.x*blockIdx.x;
    if(data[idx] == value)
        atomicMin(min_idx, idx);
}

オカレンスの数が非常に少なく、ほとんどすべてのスレッドがアトミックにアクセスしようとさえしない場合、これは実際にはそれほど悪い解決策ではないかもしれません。それ以外の場合 (検索された要素がそれほど珍しくない場合)、ワープ内の発散がはるかに多くなり、さらに悪いことに、アトミック操作が競合する可能性がはるかに高くなります。

---

編集:より洗練されたアプローチ (ただし、まだ最善ではないかもしれません) の場合は、前のステップで、入力配列の要素がそのインデックスで検索された要素と等しい場合に、インデックスの値が設定されたint配列を作成することもできます。しません:idxidxINT_MAX

indices[idx] = (data[idx]==value) ? idx : INT_MAX;

次に、そのインデックス配列に対して「古典的な」最小削減を実行して、最初に一致するインデックスを取得します。