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

Question

Я ищу стратегию оптимизации для моей программы cuda. На каждой итерации внутри цикла for моего ядра каждый поток производит оценку. Я поддерживаю общую очередь приоритетов оценок, чтобы поддерживать top-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 в этом случае будет очень скудным.

Answer 1

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

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