DeviceRadixSort::SortKeys<bool> yields compiler errors

[jaredhoberock]

The following program:

#include <cub/cub.cuh>

int main()
{
  cub::DeviceRadixSort::SortKeys<bool>;
  return 0;
}

does not compile:

$nvcc -I. sort_bool.cu
sort_bool.cu(5): warning: expression has no effect

./cub/device/dispatch/../../block_sweep/block_radix_sort_upsweep.cuh(92): error: class "cub::Traits<__nv_bool>" has no member "UnsignedBits"
          detected during:
            instantiation of class "cub::BlockRadixSortUpsweep<BlockRadixSortUpsweepPolicy, Key, Offset> [with BlockRadixSortUpsweepPolicy=cub::DeviceRadixSortDispatch<false, __nv_bool, cub::NullType, int>::PtxUpsweepPolicy, Key=__nv_bool, Offset=int]" 
./cub/device/dispatch/device_radix_sort_dispatch.cuh(80): here
            instantiation of "void cub::DeviceRadixSortUpsweepKernel<BlockRadixSortUpsweepPolicy,DESCENDING,Key,Offset>(Key *, Offset *, Offset, int, int, __nv_bool, cub::GridEvenShare<Offset>) [with BlockRadixSortUpsweepPolicy=cub::DeviceRadixSortDispatch<false, __nv_bool, cub::NullType, int>::PtxUpsweepPolicy, DESCENDING=false, Key=__nv_bool, Offset=int]" 
./cub/device/dispatch/device_radix_sort_dispatch.cuh(931): here
            instantiation of "cudaError_t cub::DeviceRadixSortDispatch<DESCENDING, Key, Value, Offset>::Dispatch(void *, size_t &, cub::DoubleBuffer<Key> &, cub::DoubleBuffer<Value> &, Offset, int, int, cudaStream_t, __nv_bool) [with DESCENDING=false, Key=__nv_bool, Value=cub::NullType, Offset=int]" 
./cub/device/device_radix_sort.cuh(331): here
            instantiation of "cudaError_t cub::DeviceRadixSort::SortKeys(void *, size_t &, cub::DoubleBuffer<Key> &, int, int, int, cudaStream_t, __nv_bool) [with Key=__nv_bool]" 

...

I tried to find out if this was intended to be supported or not. The documentation implies it is:

DeviceRadixSort can sort all of the built-in C++ numeric primitive types, e.g.: unsigned char, int,
double, etc. Although the direct radix sorting method can only be applied to unsigned integral types,
BlockRadixSort is able to sort signed and floating-point types via simple bit-wise transformations
that ensure lexicographic key ordering.

bool is indeed an integral type, so I think it's worth supporting for completeness.

[dumerrill]

Fixed in v1.5