cuda 内核中的流压缩以维护优先级队列答案

【问题标题】：Stream compaction within cuda kernel for maintaining priority queuecuda 内核中的流压缩以维护优先级队列
【发布时间】：2014-03-07 14:32:26
【问题描述】：

我正在为我的 cuda 程序寻找优化策略。在我内核的 for 循环内的每次迭代中，每个线程都会产生一个分数。我正在维护分数的共享优先级队列，以维护每个块中的前 k 个分数。请看下面的伪代码：

__global__ gpuCompute(... arguments)
{
    __shared__ myPriorityQueue[k];  //To maintain top k scores ( k < #threads in this block)
    __shared__ scoresToInsertInQueue[#threadsInBlock];
    __shared__ counter;
    for(...)       //About 1-100 million iterations
    {
        int score = calculate_score(...);
        if(score > Minimum element of P. Queue && ...)
        {
            ATOMIC Operation : localCounter = counter++;
            scoresToInsertInQueue[localCounter] = score;
        }
        __syncthreads();
        //Merge scores from scoresToInsertInQueue to myPriorityQueue
        while(counter>0)
        {
            //Parallel insertion of scoresToInsertInQueue[counter] to myPriorityQueue using the participation of k threads in this block
            counter--;  
            __syncthreads(); 
        }
        __syncthreads();
    }
}

希望上面的代码对你们有意义。现在，我正在寻找一种方法来消除原子操作开销 s.t.每个线程根据值是否应该进入优先级队列保存“1”或“0”。我想知道内核中是否有任何流压缩的实现，以便我可以将“1000000000100000000”减少到“11000000000000000000”缓冲区（或知道“1”的索引），最后在队列中插入与“1”相对应的分数。
请注意，在这种情况下，'1' 会非常稀疏。

【问题讨论】：

thrust 有 stream-compaction functions 但似乎需要您将内核分解成碎片。
@Robert：是的，我读到了这个，这个解决方案适用于这个question。将其分解为多个内核会导致性能非常差。

标签： cuda parallel-processing atomic priority-queue stream-compaction

【解决方案1】：

如果这些非常稀疏，atomic 方法可能是最快的。然而，我在这里描述的方法将具有更可预测和有界的最坏情况性能。

要在决策数组中很好地混合 1 和 0，使用并行扫描或 prefix-sum 从决策数组中构建插入点索引数组可能会更快：

假设我有一个阈值决策，选择分数 > 30 进入队列。我的数据可能如下所示：

scores:     30  32  28  77  55  12  19
score > 30:  0   1   0   1   1   0   0
insert_pt:   0   0   1   1   2   3   3    (an "exclusive prefix sum")

然后每个线程做出如下存储选择：

if (score[threadIdx.x] > 30) temp[threadIdx.x] = 1;
else temp[threadIdx.x] = 0;
__syncthreads();
// perform exclusive scan on temp array into insert_pt array
__syncthreads();
if (temp[threadIdx.x] == 1)
  myPriorityQueue[insert_pt[threadIdx.x]] = score[threadIdx.x];

CUB 具有快速并行前缀扫描。

【讨论】：