[BUG]: Improve error reporting for segmented reduce

[gevtushenko]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this bug and that I agree to the Code of Conduct

Type of Bug

Compile-time Error

Component

CUB

Describe the bug

As of #414 we use common value type of offset begin / end iterators as offset type in segmented reduce. Incorrect use of API is hard to diagnose. For instance, user can pass offset iterator with float as its value type. Float is not expected on the implementation side of segmented reduce, leading to many errors that are hard to associate with API misuse.

How to Reproduce

#include <cub/cub.cuh>

int main() {
  using offset_t = float;  // error
  // using offset_t = int; // ok
  float *d_in{}, *d_out{};
  offset_t *d_begin_offsets{}, *d_end_offsets{};
  std::size_t temp_storage_bytes{};
  std::uint8_t *d_temp_storage{};
  cub::DeviceSegmentedReduce::Sum(d_temp_storage, temp_storage_bytes, d_in, d_out, 0, d_begin_offsets, d_end_offsets);
}

Expected behavior

We should use combination of SFINAE and static asserts like we do in device radix sort to improve error reporting.

Reproduction link

No response

Operating System

No response

nvidia-smi output

No response

NVCC version

No response