  *
  * Dispatch to shared atomic entrypoint
  */
-template <int BINS, int CHANNELS, int ACTIVE_CHANNELS, typename InputIterator, typename HistoCounter>
+template <int BINS, int CHANNELS, int ACTIVE_CHANNELS, typename SampleT, typename InputIterator, typename HistoCounter>
 __host__ __device__ __forceinline__
 cudaError_t Dispatch(
     Int2Type<DEVICE_HISTO_SHARED_ATOMIC> algorithm,


    printf("%s cub::DeviceHistogram %s %d %s samples (%dB), %d bins, %d channels, %d active channels, gen-mode %s\n",
        (CDP) ? "CDP device invoked" : "Host-invoked",
        (ALGORITHM == DEVICE_HISTO_SHARED_ATOMIC) ? "satomic" : (ALGORITHM == DEVICE_HISTO_GLOBAL_ATOMIC) ? "gatomic" : "sort",
        num_samples,
        type_string,
        (int) sizeof(SampleT),


    void                *d_temp_storage,
    size_t              &temp_storage_bytes,
    SampleT             *d_samples,
    InputIterator       d_sample_itr,
    HistoCounter        *d_histograms[ACTIVE_CHANNELS],
    int                 num_samples,
    cudaStream_t        stream,
