Fix DeviceHistogram bug: non-native input iterators take the correct code path.
Fix DeviceHistogram bug: non-native input iterators take the correct code path.

Previously, when passing in an, e.g. CountingInputIterator, the kernel
would fail with the error:

======== Invalid __global__ read of size 4
=========     by thread (0,0,0) in block (0,0,0)
=========     Address 0x00000000 is out of bounds

because the wrong branch tried to read from d_native_samples (which is null).

NVIDIA/cub,NVlabs/CUB,NVlabs/cub