Sorting (small) arrays by key in CUDA

cub, cuda, parallel-processing, reduce, sorting

Solution

You can use thrust to do this.

Use thrust::sort_by_key followed by thrust::reduce_by_key

Here's an example:

#include <iostream>
#include <thrust/device_vector.h>
#include <thrust/copy.h>
#include <thrust/sort.h>
#include <thrust/reduce.h>
#include <thrust/sequence.h>

#define N 12
typedef thrust::device_vector<int>::iterator dintiter;
int main(){

  thrust::device_vector<int> keys(N);
  thrust::device_vector<int> values(N);
  thrust::device_vector<int> new_keys(N);
  thrust::device_vector<int> new_values(N);
  thrust::sequence(keys.begin(), keys.end());
  thrust::sequence(values.begin(), values.end());

  keys[3] = 1;
  keys[9] = 1;
  keys[8] = 2;
  keys[7] = 4;

  thrust::sort_by_key(keys.begin(), keys.end(), values.begin());
  thrust::pair<dintiter, dintiter> new_end;
  new_end = thrust::reduce_by_key(keys.begin(), keys.end(), values.begin(), new_keys.begin(), new_values.begin());

  std::cout << "results  values:" << std::endl;
  thrust::copy(new_values.begin(), new_end.second, std::ostream_iterator<int>( std::cout, " "));
  std::cout << std::endl << "results keys:" << std::endl;
  thrust::copy(new_keys.begin(), new_end.first, std::ostream_iterator<int>( std::cout, " "));
  std::cout << std::endl;

  return 0;
}

Problem

I'm trying to write a function that takes a block of unsorted key/value pairs such as ``` <7, 4> <2, 8> <3, 1> <2, 2> <1, 5> <7, 1> <3, 8> <7, 2> ``` and sorts them by key while reducing the values of pairs with the same key: ``` <1, 5> <2, 10> <3, 9> <7, 7> ``` Currently, I'm using a `__device__` function like the one below which is essentially a bitonic sort that will combine values of the same key and set the old data to an infinitely large value (just using `99` for now) so that a subsequent bitonic sort will sift them to the bottom and the array cut by the value of `int *` removed. ``` __device__ void interBitonicSortReduce(int2 *sdata, int tid, int recordNum, int *removed) { int n = MIN(DEFAULT_DIMBLOCK, recordNum); for (int k = 2; k <= n; k *= 2) { for (int j = k / 2; j > 0; j /= 2) { int ixj = tid ^ j; if (ixj > tid) { if (sdata[tid].x == sdata[ixj].x && sdata[tid].x < 99) { atomicAdd(&sdata[tid].y, sdata[ixj].y); sdata[ixj].x = 99; sdata[ixj].y = 99; atomicAdd(removed, 1); } if ((tid & k) == 0 && sdata[tid].x > sdata[ixj].x) swapData2(sdata[tid], sdata[ixj]); if ((tid & k) != 0 && sdata[tid].x < sdata[ixj].x) swapData2(sdata[tid], sdata[ixj]); __syncthreads(); } } } } ``` This works just fine for small sets of data but with larger sets (though still within the size of a single block) a single call just won't do it. Is it wise to try to combine the sorting and the reduction in the same function? Obviously the function would need to be called more than once but is it possible to determine exactly how many times it needs to be called to exhaust all the data based on its size? Or should I preform the reduction separately with something like this: ``` __device__ int interReduce(int2 *sdata, int tid) { int index = tid; while (sdata[index].x == sdata[tid].x) { index--; if (index < 0) break; } if (index+1 != tid) { atomicAdd(&sdata[index+1].y, sdata[tid].y); sdata[tid].x = 99; sdata[tid].y = 99; return 1; } return 0; } ``` I'm trying to come up with the most efficient solution, but my experience with CUDA and parallel algorithms is limited.

Original source