1

我正在为我的 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();
    }
}

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

4

1 回答 1

2

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

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

假设我有一个阈值决定，选择分数 > 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具有快速的并行前缀扫描。

于 2014-02-10T19:51:39.513 回答