データがキューに入れられていない配列にデータを効率的に挿入する

Question

私はCuda内に2つの配列を持っています。

int *main; // unsorted 
int *source; // sorted 

私のアルゴリズムの一部は、ソース配列からメイン配列に新しいデータを定期的に挿入する必要があります。メイン配列内の位置がゼロの場合は、空であるとみなされるため、ソース配列の値を取り込むことができます。

これを実行する最も効率的な方法が何であるか、私はいくつかのアプローチを試しましたが、ここで行われるパフォーマンスの向上はまだあると思います。

現在、メイン配列の内容をメイン配列の最後までシャッフルして、配列の先頭にすべてのゼロ値を残して、挿入を行いますソースからはほんのわずかです。ソートは、32ビットではなく1つのビットを反復するように変更されました。これは、入力上の単純なスイッチで動作します。

input[i] = source[i] > 1 ? 1 : 0 

これは既に効率的な方法ですか？私は戦術的に配備されたatomicAddを使って何かを得ることができないのだろうかと思っています。

__global__ void find(int *destination, int *indices, const int N) 
{ 
    int idx = blockIdx.x * blockDim.x + threadIdx.x; 

    if((destination[idx] == 0)&&(count<elements_to_add)) 
    { 
     indices[count] = idx; 
     atomicAdd(&count, 1); 
    } 
} 

__global__ void insert(int *destination, int *indices, int *source, const int N) 
{ 
    int idx = blockIdx.x * blockDim.x + threadIdx.x; 

    if((source[idx] > 0)&&(indices[idx] > 0)) 
    { 
     destination[indices[idx]] = source[idx]; 
    } 
} 

find<<<G,T>>>(...); 
insert<<<G,T>>>(...); 

現時点ではソース配列を介して多くのアイテムを挿入していませんが、将来変更される可能性があります。

これは、以前に解決された共通の問題であるように思えます。推力ライブラリが役立つかもしれないと思っていますが、私が試しているものに対しては、達成するために（私はすでに持っているコードとあまりよくきちんと合っていません）

経験豊富なCuda開発者の意見が高く評価されました！

Answer 1

ストリームコンパクション手順に分類される検索アルゴリズムと、散布手順に分類される挿入を分離することができます。ただし、2つの機能をマージすることはできます。

と仮定すると、srcPtrは、その内容がグローバルメモリ内に存在し、カーネルが起動する前に既にゼロに設定されていることを前提としています。

__global__ void find_and_insert(int* destination, int const* source, int const N, int* srcPtr) { // Assuming N is the length of the destination buffer and also the length of the source buffer is less than N. 

int const idx = blockIdx.x * blockDim.x + threadIdx.x; 

// Get the assigned element. 
int const dstElem = destination[ idx ]; 
bool const pred = (dstElem == 0); 

// Intra-warp binary reduction to count the total number of lanes with empty elements. 
int const predBallot = __ballot(pred); 
int const intraWarpRed = __popc(predBallot); 

// Warp-aggregated atomics to reduce the contention over the srcPtr content. 
unsigned int laneID; asm("mov.u32 %0, %laneid;" : "=r"(laneID)); //const uint laneID = tidWithinCTA & (WARP_SIZE - 1); 
int posW; 
if(laneID == 0) 
    posW = atomicAdd(srcPtr, intraWarpRed); 
posW = __shfl(posW, 0); 

// Threads that have found empty elements can fill out their assigned positions from the src. Intra-warp binary prefix sum is used here. 
uint laneMask; asm("mov.u32 %0, %lanemask_lt;" : "=r"(laneMask)); //const uint laneMask = 0xFFFFFFFF >> (WARP_SIZE - laneID) ; 
int const positionToRead = posW + __popc(predBallot & laneMask); 
if(pred) 
    destination[ idx ] = source[ positionToRead ]; 

} 

いくつかのこと：

1. このカーネルは、あなたがそれを行うことができますどのように単なる提案です。ここで、ワープの内部のスレッドは、そのタスクに関して協力します。バイナリ縮小と接頭辞の合計をスレッドブロックに拡張することができます。
2. 私はこのカーネルをブラウザ内に書きましたが、テストしていません。ので注意してください。
3. デザイン全体が新しいものではありません。同様のアプローチが実施されており（例えば、this paper）、ほとんどがthe work done by Mark Harris and Michael Garlandに基づいています。