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).

cub/agent/agent_histogram.cuh

@@ -746,7 +746,7 @@ struct AgentHistogram
                                         ((row_bytes & pixel_mask) == 0);                        // number of row-samples is a multiple of the alignment of the pixel
 
         // Whether rows are aligned and can be vectorized
-        if (quad_aligned_rows || pixel_aligned_rows)
+        if ((d_native_samples != nullptr) && (quad_aligned_rows || pixel_aligned_rows))
             ConsumeTiles<true>(num_row_pixels, num_rows, row_stride_samples, tiles_per_row, tile_queue, Int2Type<IS_WORK_STEALING>());
         else
             ConsumeTiles<false>(num_row_pixels, num_rows, row_stride_samples, tiles_per_row, tile_queue, Int2Type<IS_WORK_STEALING>());