CUDA Standard Algorithms » Parallel Sort

Sort a Range of Items

tf::cuda_sort performs an in-place parallel sort over a range of elements specified by [first, last) using the given comparator. The following code sorts one million random integers in an increasing order on a GPU.

const size_t N = 1000000;
int* vec = tf::cuda_malloc_shared<int>(N);  // vector

// initializes the data
for(size_t i=0; i<N; i++) 
  vec[i] = rand();
} 
tf::cuda_sort(tf::cudaDefaultExecutionPolicy{}, vec, vec+N);

assert(std::is_sorted(vec, vec+N));

You can specify a comparator to tf::cuda_sort to alter the sorting order. For example, the following code sorts one million random integers in an decreasing order on a GPU.

const size_t N = 1000000;
int* vec = tf::cuda_malloc_shared<int>(N);  // vector

// initializes the data
for(size_t i=0; i<N; i++) 
  vec[i] = rand();
} 
tf::cuda_sort(tf::cudaDefaultExecutionPolicy{}, 
  vec, vec+N, [] __device__ (int a, int b) { return a > b; }
);

assert(std::is_sorted(vec, vec+N, [](int a, int b){ return a > b; }));

Invoke Sort Kernels Asynchronously

By default, tf::cuda_sort blocks until all kernels finish. You can use tf::cuda_sort_async to invoke these kernels asynchronously and explicitly synchronize them at another place of your program. Since our sorting kernels rely on additional device memory, you need to provide a buffer of size in bytes equal to (or larger than) the value returned by tf::cuda_sort_buffer_size.

const size_t N = 100000;
int* vec = tf::cuda_malloc_shared<int>(N);  // vector

// initializes the data
for(size_t i=0; i<N; i++) 
  vec[i] = rand();
} 

// queries the required buffer size to sort N elements using the default policy
auto bytes = tf::cuda_sort_buffer_size<tf::cudaDefaultExecutionPolicy, int>(N);
void* buffer = tf::cuda_malloc_device<std::byte>(bytes);

// invoke the sort kernels asynchronously
tf::cuda_sort_async(tf::cudaDefaultExecutionPolicy{my_stream},
  vec, vec+N, [] __device__ (int a, int b) { return a < b; }, buffer
);
// here, vec may not be sorted yet, as kernels are still running ...

// synchronize the stream
cudaStreamSynchronize(my_stream);

assert(std::is_sorted(vec, ved+N));

The allocated buffer must remain alive until the asynchronous call completes.