Device-scope segmented reduce temp storage allocation issue

[gevtushenko]

Device-scope segmented reduce contains the following short-circuit:

if (num_segments <= 0)
{
  return cudaSuccess;
}

which happens before temp_storage_bytes is initialized. The issue might take place:

size_t temp_storage_bytes; // <- uninitialized
cub::DeviceSegmentedReduce::Reduce(nullptr, temp_storage_bytes);
thrust::device_vector<char> temp_storage(temp_storage_bytes); // <- allocate random amount of memory

We guarantee to overwrite temp_storage_bytes in the docs. Therefore, we should fix this.