CUDA Standard Algorithms » Parallel Merge

Merge Two Sorted Ranges of Items

tf::cuda_merge merges two sorted ranges of items into a sorted range. The following code merges two sorted arrays input_1 and input_2, each of 1000 items, into a sorted array output of 2000 items.

const size_t N = 1000;
int* input_1 = tf::cuda_malloc_shared<int>(N);    // input vector 1
int* input_2 = tf::cuda_malloc_shared<int>(N);    // input vector 2
int* output  = tf::cuda_malloc_shared<int>(2*N);  // output vector

// initializes the data
for(size_t i=0; i<N; i++) {
  input_1[i] = rand()%100;
  input_2[i] = rand()%100;
}
std::sort(input_1, input1 + N);
std::sort(input_2, input2 + N);

// merge input_1 and input_2 to output
tf::cuda_merge(tf::cudaDefaultExecutionPolicy{}, 
  input_1, input_1 + N, input_2, input_2 + N, output, 
  []__device__ (int a, int b) { return a < b; }  // comparator
);

Invoke Merge Kernels Asynchronously

By default, tf::cuda_merge blocks until all kernels finish. You can invoke these kernels asynchronously and explicitly synchronize them at another place of your program. Since our merge 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_merge_buffer_size.

const size_t N = 1000;
int* input_1 = tf::cuda_malloc_shared<int>(N);    // input vector 1
int* input_2 = tf::cuda_malloc_shared<int>(N);    // input vector 2
int* output  = tf::cuda_malloc_shared<int>(2*N);  // output vector

// initializes the data
for(size_t i=0; i<N; i++) {
  input_1[i] = rand()%100;
  input_2[i] = rand()%100;
}
std::sort(input_1, input1 + N);
std::sort(input_2, input2 + N);

// queries the required buffer size to merge two arrays using the given policy
auto bytes = tf::cuda_merge_buffer_size<tf::cudaDefaultExecutionPolicy>(N, N);
void* buffer = tf::cuda_malloc_device<std::byte>(bytes);

// invoke the merge kernels asynchronously
tf::cuda_merge_async(tf::cudaDefaultExecutionPolicy{my_stream},
  input_1, input_1 + N, input_2, input_2 + N, output, 
  [] __device__ (int a, int b) { return a < b; }, buffer
);
// here, output may not be ready yet, as kernels are still running ...

// synchronize the stream to make output ready
cudaStreamSynchronize(my_stream);

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