Floating point exception (core dumped) during cub::DeviceRadixSort::SortKeys

[jaredhoberock]

The following program:

#include <thrust/device_vector.h>
#include <thrust/copy.h>
#include <thrust/system/cuda/execution_policy.h>
#include <cub/device/device_radix_sort.cuh>


template<typename T>
void cub_sort(T* first, size_t n)
{
  cudaError_t error = cudaSuccess;

  cub::DoubleBuffer<T> double_buffer;

  // measure the number of additional temporary storage bytes required
  size_t num_additional_temp_storage_bytes = 0;
  error = cub::DeviceRadixSort::SortKeys(0, num_additional_temp_storage_bytes, double_buffer, n);
  if(error)
  {
    throw thrust::system_error(error, thrust::cuda_category(), "after cub::DeviceRadixSort::SortKeys(0)");
  }

  // allocate temporary storage for double buffer as well as additional bytes
  // XXX seems like we should align up the additional bytes
  size_t num_double_buffer_bytes = n * sizeof(T);
  thrust::device_vector<char> temporary_storage(num_double_buffer_bytes + num_additional_temp_storage_bytes);

  void* additional_temp_storage_ptr = thrust::raw_pointer_cast(temporary_storage.data() + num_double_buffer_bytes);

  double_buffer.d_buffers[0] = thrust::raw_pointer_cast(&*first);
  double_buffer.d_buffers[1] = reinterpret_cast<T*>(thrust::raw_pointer_cast(temporary_storage.data()));

  error = cub::DeviceRadixSort::SortKeys(additional_temp_storage_ptr,
                                         num_additional_temp_storage_bytes,
                                         double_buffer,
                                         static_cast<int>(n));

  if(error)
  {
    throw thrust::system_error(error, thrust::cuda_category(), "after cub::DeviceRadixSort::SortKeys(1)");
  }

  if(double_buffer.Current() != 0)
  {
    T* temp_ptr = reinterpret_cast<T*>(double_buffer.d_buffers[1]);
    thrust::copy(thrust::cuda::par, temp_ptr, temp_ptr + n, first);
  }
}


int main()
{
   thrust::device_vector<int> data(0);

   cub_sort(thrust::raw_pointer_cast(data.data()), data.size());

  return 0;
}

Generates the following output at runtime:

$ nvcc -arch=sm_35 -I. cub_repro.cu -run
Floating point exception (core dumped)

Presumably a divide by zero occurs due to n == 0.

[dumerrill]

Fixed in v1.5