c++ - すべてのスレッドではなく、CUDA カーネルで配列またはリストを埋める

Question

基本的に、カーネルに if() があり、条件が検証された場合は、動的リストまたは配列に新しい値を格納したいと考えています。問題は、すべてのカーネルに入力されるわけではないため、threadIdx を使用できないことです。

何かのようなもの：

__global__ void myKernel(customType *c)
{
    int i = threadIdx.x;
    //whatever
    if(condition)
        c->pop(newvalue)
}

実際、ac[i]=newvalue は避けたいと思います。なぜなら、ホストコードに for ループを使用して値が挿入されているかどうかを最後にすべての c[i] で確認し、別の構造を適切に埋める必要があるからです。スラストについて考えましたが、私の「単純な」問題にはやり過ぎのようです。

回避策を見つけるのを手伝っていただければ幸いです。

score 4 · Accepted Answer

私の理解が正しければ、ストリーム圧縮について説明しています。すべてのスレッドではなく一部のスレッドが値を作成し、それらの値をギャップなしで配列に格納する必要があります。

これを実装する 1 つの方法は、 Thrustで利用可能なストリーム圧縮アルゴリズムを使用することです(この例を参照してください)。これには、2 つのパスで操作を実行する必要があることに注意してください。

(グリッド全体ではなく) 単一のスレッドブロック内からこれを実行している場合は、CUBを確認することもできます。各スレッドは、値を保存するかどうかを示すフラグを計算し、フラグに対してプレフィックスサムを実行して、リスト内の各スレッドのオフセットを決定してから、保存を行います。

2 に答える 2