Сжатие потока в ядре cuda для поддержки очереди с приоритетами

Question

Я ищу стратегию оптимизации для моей программы 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();
    }
}

Надеюсь, что приведенный выше код имеет смысл для вас, ребята. Теперь я ищу способ удалить накладные расходы атомарной операции, когда каждый поток сохраняет "1" или "0" в зависимости от того, должно ли значение идти в приоритетную очередь или нет. Мне интересно, есть ли какая-либо реализация потокового сжатия в ядре, чтобы я мог уменьшить буфер "1000000000100000000" до "11000000000000000000" (или знать индекс "1") и, наконец, вставить оценки, соответствующие "1", в очередь.
Обратите внимание, что "1" будет очень редким в этом случае.

Answer 1

Если они очень редки, atomic метод может быть самым быстрым. Однако описанный здесь метод будет иметь более предсказуемую и ограниченную производительность в худшем случае.

Для хорошего сочетания единиц и нулей в вашем массиве решений может быть быстрее использовать параллельное сканирование или префиксную сумму для построения массива индекса точки вставки из массива решений:

Предположим, у меня есть решение о пороговом значении, которое выбирает баллы> 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 имеет быстрое параллельное сканирование префиксов.