DeviceHistogram fails to compile trying to use atomicAdd on unsupported types

I am trying to compile the following CUB test program -

```
#include

int main()
{
uint32_t* d_samples;
uint8_t* d_histogram;
uint8_t* d_levels;
size_t temp_storage_bytes;
int num_levels;
size_t num_samples;
cub::DeviceHistogram::HistogramRange(nullptr,
temp_storage_bytes,
d_samples,
d_histogram,
num_levels,
d_levels,
num_samples);
return 0;
}
```

using the following command line -
`/usr/local/cuda-11.0/bin/nvcc cubDeviceHistogramTest.cu -c -o cubDeviceHistogramTest.o -ccbin /usr/bin/cc -m64 -Xcompiler ,\"-Wall\",\"-Wextra\",\"-Werror\",\"-fPIC\",\"-g\" -std=c++14 --expt-extended-lambda --expt-relaxed-constexpr -lineinfo --generate-code arch=compute_61,code=sm_61 --generate-code arch=compute_70,code=sm_70 -DNVCC -I/usr/local/cuda-11.0/include`

but I see the following error -
```
/usr/local/cuda-11.0/include/cub/device/dispatch/../../agent/agent_histogram.cuh(314): error: no instance of overloaded function "atomi
cAdd" matches the argument list
argument types are: (uint8_t *, uint8_t)
detected during:
instantiation of "void cub::AgentHistogram::StoreOutput(CounterT **) [with AgentHistogramPolic
yT=cub::DipatchHistogram<1, 1, uint32_t *, uint8_t, uint8_t, int>::PtxHistogramSweepPolicy, PRIVATIZED_SMEM_BINS=256, NUM_CHANNELS=1, N
UM_ACTIVE_CHANNELS=1, SampleIteratorT=uint32_t *, CounterT=uint8_t, PrivatizedDecodeOpT=cub::DipatchHistogram<1, 1, uint32_t *, uint8_t
, uint8_t, int>::SearchTransform, OutputDecodeOpT=cub::DipatchHistogram<1, 1, uint32_t *, uint8_t, uint8_t, int>::PassThruTr
ansform, OffsetT=int, PTX_ARCH=700]"
(329): here
instantiation of "void cub::AgentHistogram::StoreSmemOutput() [with AgentHistogramPolicyT=cub:
:DipatchHistogram<1, 1, uint32_t *, uint8_t, uint8_t, int>::PtxHistogramSweepPolicy, PRIVATIZED_SMEM_BINS=256, NUM_CHANNELS=1, NUM_ACTI
VE_CHANNELS=1, SampleIteratorT=uint32_t *, CounterT=uint8_t, PrivatizedDecodeOpT=cub::DipatchHistogram<1, 1, uint32_t *, uint8_t, uint8
_t, int>::SearchTransform, OutputDecodeOpT=cub::DipatchHistogram<1, 1, uint32_t *, uint8_t, uint8_t, int>::PassThruTransform
, OffsetT=int, PTX_ARCH=700]"
(774): here
instantiation of "void cub::AgentHistogram::StoreOutput() [with AgentHistogramPolicyT=cub::DipatchHistogram<1, 1, uint32_t *, uint8_t, uint8_t, int>::PtxHistogramSweepPolicy, PRIVATIZED_SMEM_BINS=256, NUM_CHANNELS=1, NUM_ACTIVE_$HANNELS=1, SampleIteratorT=uint32_t *, CounterT=uint8_t, PrivatizedDecodeOpT=cub::DipatchHistogram<1, 1, uint32_t *, uint8_t, uint8_t,
int>::SearchTransform, OutputDecodeOpT=cub::DipatchHistogram<1, 1, uint32_t *, uint8_t, uint8_t, int>::PassThruTransform, O$fsetT=int, PTX_ARCH=700]"
/usr/local/cuda-11.0/include/cub/device/dispatch/dispatch_histogram.cuh(154): here
instantiation of "void cub::DeviceHistogramSweepKernel(SampleIteratorT, cub::ArrayWrapper, cub::ArrayWrapper, cub::ArrayWrapper, cub::ArrayWrapper, cub::ArrayWrapper, cub::ArrayWrapper) [with AgentHistogramPolicyT=cub::DipatchHistogram<1, 1, uint32_t *, uint8_t, uin
t8_t, int>::PtxHistogramSweepPolicy, PRIVATIZED_SMEM_BINS=256, NUM_CHANNELS=1, NUM_ACTIVE_CHANNELS=1, SampleIteratorT=uint32_t *, Count
erT=uint8_t, PrivatizedDecodeOpT=cub::DipatchHistogram<1, 1, uint32_t *, uint8_t, uint8_t, int>::SearchTransform, OutputDeco
deOpT=cub::DipatchHistogram<1, 1, uint32_t *, uint8_t, uint8_t, int>::PassThruTransform, OffsetT=int]"
/usr/local/cuda-11.0/include/cub/device/dispatch/dispatch_histogram.cuh(792): here
instantiation of "cudaError_t cub::DipatchHistogram::DispatchRange(void *, size_t &, SampleIteratorT, CounterT **, int *, LevelT **, OffsetT, OffsetT, OffsetT, cudaStream_t, __nv_
bool, cub::Int2Type<0>) [with NUM_CHANNELS=1, NUM_ACTIVE_CHANNELS=1, SampleIteratorT=uint32_t *, CounterT=uint8_t, LevelT=uint8_t, Offs
etT=int]"
/usr/local/cuda-11.0/include/cub/device/device_histogram.cuh(849): here
instantiation of "cudaError_t cub::DeviceHistogram::MultiHistogramRange(void *, size_t &, SampleIteratorT, CounterT **, int *, LevelT **, OffsetT, OffsetT, size_t, cudaStream_t, __nv_b
ool) [with NUM_CHANNELS=1, NUM_ACTIVE_CHANNELS=1, SampleIteratorT=uint32_t *, CounterT=uint8_t, LevelT=uint8_t, OffsetT=int]"
/usr/local/cuda-11.0/include/cub/device/device_histogram.cuh(557): here
instantiation of "cudaError_t cub::DeviceHistogram::HistogramRange(void *, size_t &, SampleIteratorT, CounterT *, int, Leve
lT *, OffsetT, cudaStream_t, __nv_bool) [with SampleIteratorT=uint32_t *, CounterT=uint8_t, LevelT=uint8_t, OffsetT=int]"
cubDeviceHistogramTest.cu(14): here
```

Looks like this is a bug.